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>
diff --git a/BUILD.gn b/BUILD.gn
index 12578bb..7a14689 100644
--- a/BUILD.gn
+++ b/BUILD.gn
@@ -286,8 +286,6 @@
"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 @@
"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",
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index f79ae1a..d9a15f7 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -100,8 +100,6 @@
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 @@
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
diff --git a/src/ast/identifier_expression.h b/src/ast/identifier_expression.h
index 1021bdd..54cfdba 100644
--- a/src/ast/identifier_expression.h
+++ b/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 @@
/// @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,9 +69,7 @@
Symbol const sym_;
- Intrinsic intrinsic_ = Intrinsic::kNone; // Semantic info
- std::unique_ptr<intrinsic::Signature> intrinsic_sig_; // Semantic info
- bool is_swizzle_ = false; // Semantic info
+ bool is_swizzle_ = false; // Semantic info
};
} // namespace ast
diff --git a/src/ast/intrinsic.cc b/src/ast/intrinsic.cc
deleted file mode 100644
index 6a16abb..0000000
--- a/src/ast/intrinsic.cc
+++ /dev/null
@@ -1,294 +0,0 @@
-// Copyright 2020 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/ast/intrinsic.h"
-
-namespace tint {
-namespace ast {
-
-std::ostream& operator<<(std::ostream& out, Intrinsic i) {
- /// The emitted name matches the spelling in the WGSL spec.
- /// 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;
-TextureSignature::~TextureSignature() = default;
-
-TextureSignature::Parameters::Index::Index() = default;
-TextureSignature::Parameters::Index::Index(const Index&) = default;
-
-bool IsCoarseDerivative(Intrinsic i) {
- return i == Intrinsic::kDpdxCoarse || i == Intrinsic::kDpdyCoarse ||
- i == Intrinsic::kFwidthCoarse;
-}
-
-bool IsFineDerivative(Intrinsic i) {
- return i == Intrinsic::kDpdxFine || i == Intrinsic::kDpdyFine ||
- i == Intrinsic::kFwidthFine;
-}
-
-bool IsDerivative(Intrinsic i) {
- return i == Intrinsic::kDpdx || i == Intrinsic::kDpdy ||
- i == Intrinsic::kFwidth || IsCoarseDerivative(i) ||
- IsFineDerivative(i);
-}
-
-bool IsFloatClassificationIntrinsic(Intrinsic i) {
- return i == Intrinsic::kIsFinite || i == Intrinsic::kIsInf ||
- i == Intrinsic::kIsNan || i == Intrinsic::kIsNormal;
-}
-
-bool IsTextureIntrinsic(Intrinsic i) {
- return IsImageQueryIntrinsic(i) || i == Intrinsic::kTextureLoad ||
- i == Intrinsic::kTextureSample ||
- i == Intrinsic::kTextureSampleLevel ||
- i == Intrinsic::kTextureSampleBias ||
- i == Intrinsic::kTextureSampleCompare ||
- i == Intrinsic::kTextureSampleGrad || i == Intrinsic::kTextureStore;
-}
-
-bool IsImageQueryIntrinsic(Intrinsic i) {
- return i == ast::Intrinsic::kTextureDimensions ||
- i == Intrinsic::kTextureNumLayers ||
- i == Intrinsic::kTextureNumLevels ||
- i == Intrinsic::kTextureNumSamples;
-}
-
-} // namespace intrinsic
-} // namespace ast
-} // namespace tint
diff --git a/src/ast/intrinsic.h b/src/ast/intrinsic.h
deleted file mode 100644
index 90477bc..0000000
--- a/src/ast/intrinsic.h
+++ /dev/null
@@ -1,202 +0,0 @@
-// Copyright 2020 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_AST_INTRINSIC_H_
-#define SRC_AST_INTRINSIC_H_
-
-#include <ostream>
-
-namespace tint {
-namespace ast {
-
-enum class Intrinsic {
- kNone = -1,
-
- kAbs,
- kAcos,
- kAll,
- kAny,
- kArrayLength,
- kAsin,
- kAtan,
- kAtan2,
- kCeil,
- kClamp,
- kCos,
- kCosh,
- kCountOneBits,
- kCross,
- kDeterminant,
- kDistance,
- kDot,
- kDpdx,
- kDpdxCoarse,
- kDpdxFine,
- kDpdy,
- kDpdyCoarse,
- kDpdyFine,
- kExp,
- kExp2,
- kFaceForward,
- kFloor,
- kFma,
- kFract,
- kFrexp,
- kFwidth,
- kFwidthCoarse,
- kFwidthFine,
- kInverseSqrt,
- kIsFinite,
- kIsInf,
- kIsNan,
- kIsNormal,
- kLdexp,
- kLength,
- kLog,
- kLog2,
- kMax,
- kMin,
- kMix,
- kModf,
- kNormalize,
- kPow,
- kReflect,
- kReverseBits,
- kRound,
- kSelect,
- kSign,
- kSin,
- kSinh,
- kSmoothStep,
- kSqrt,
- kStep,
- kTan,
- kTanh,
- kTextureDimensions,
- kTextureLoad,
- kTextureNumLayers,
- kTextureNumLevels,
- kTextureNumSamples,
- kTextureSample,
- kTextureSampleBias,
- kTextureSampleCompare,
- kTextureSampleGrad,
- kTextureSampleLevel,
- kTextureStore,
- kTrunc
-};
-
-/// Emits the name of the intrinsic function. The spelling,
-/// including case, matches the name in the WGSL spec.
-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.
-bool IsCoarseDerivative(Intrinsic i);
-
-/// Determines if the given `i` is a fine derivative
-/// @param i the intrinsic
-/// @returns true if the given derivative is fine.
-bool IsFineDerivative(Intrinsic i);
-
-/// Determine if the given `i` is a derivative intrinsic
-/// @param i the intrinsic
-/// @returns true if the given `i` is a derivative intrinsic
-bool IsDerivative(Intrinsic i);
-
-/// Determines if the given `i` is a float classification intrinsic
-/// @param i the intrinsic
-/// @returns true if the given `i` is a float intrinsic
-bool IsFloatClassificationIntrinsic(Intrinsic i);
-
-/// Determines if the given `i` is a texture operation intrinsic
-/// @param i the intrinsic
-/// @returns true if the given `i` is a texture operation intrinsic
-bool IsTextureIntrinsic(Intrinsic i);
-
-/// Determines if the given `i` is a image query intrinsic
-/// @param i the intrinsic
-/// @returns true if the given `i` is a image query intrinsic
-bool IsImageQueryIntrinsic(Intrinsic i);
-
-} // namespace intrinsic
-} // namespace ast
-} // namespace tint
-
-#endif // SRC_AST_INTRINSIC_H_
diff --git a/src/program.cc b/src/program.cc
index 269c738..09447aa 100644
--- a/src/program.cc
+++ b/src/program.cc
@@ -104,7 +104,7 @@
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;
}
diff --git a/src/program.h b/src/program.h
index 8954b7f..e872f36 100644
--- a/src/program.h
+++ b/src/program.h
@@ -119,7 +119,7 @@
/// @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
diff --git a/src/reader/spirv/function.cc b/src/reader/spirv/function.cc
index fa4300d..1bb6724 100644
--- a/src/reader/spirv/function.cc
+++ b/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 @@
return "";
}
-// Returns the WGSL standard library function instrinsic for the
-// given instruction, or ast::Intrinsic::kNone
-ast::Intrinsic GetIntrinsic(SpvOp opcode) {
+// Returns the WGSL standard library function intrinsic for the
+// given instruction, or semantic::Intrinsic::kNone
+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 @@
}
const auto intrinsic = GetIntrinsic(opcode);
- if (intrinsic != ast::Intrinsic::kNone) {
+ if (intrinsic != semantic::Intrinsic::kNone) {
return MakeIntrinsicCall(inst);
}
@@ -3990,7 +3990,6 @@
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 @@
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 =
diff --git a/src/semantic/call.h b/src/semantic/call.h
new file mode 100644
index 0000000..c1459d5
--- /dev/null
+++ b/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_
diff --git a/src/semantic/intrinsic.cc b/src/semantic/intrinsic.cc
new file mode 100644
index 0000000..f09063c
--- /dev/null
+++ b/src/semantic/intrinsic.cc
@@ -0,0 +1,220 @@
+// Copyright 2020 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/intrinsic.h"
+
+namespace tint {
+namespace semantic {
+
+std::ostream& operator<<(std::ostream& out, Intrinsic i) {
+ out << intrinsic::str(i);
+ return out;
+}
+
+namespace intrinsic {
+
+const char* str(Intrinsic i) {
+ /// The emitted name matches the spelling in the WGSL spec.
+ /// including case.
+ switch (i) {
+ 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 ||
+ i == Intrinsic::kFwidthCoarse;
+}
+
+bool IsFineDerivative(Intrinsic i) {
+ return i == Intrinsic::kDpdxFine || i == Intrinsic::kDpdyFine ||
+ i == Intrinsic::kFwidthFine;
+}
+
+bool IsDerivative(Intrinsic i) {
+ return i == Intrinsic::kDpdx || i == Intrinsic::kDpdy ||
+ i == Intrinsic::kFwidth || IsCoarseDerivative(i) ||
+ IsFineDerivative(i);
+}
+
+bool IsFloatClassificationIntrinsic(Intrinsic i) {
+ return i == Intrinsic::kIsFinite || i == Intrinsic::kIsInf ||
+ i == Intrinsic::kIsNan || i == Intrinsic::kIsNormal;
+}
+
+bool IsTextureIntrinsic(Intrinsic i) {
+ return IsImageQueryIntrinsic(i) || i == Intrinsic::kTextureLoad ||
+ i == Intrinsic::kTextureSample ||
+ i == Intrinsic::kTextureSampleLevel ||
+ i == Intrinsic::kTextureSampleBias ||
+ i == Intrinsic::kTextureSampleCompare ||
+ i == Intrinsic::kTextureSampleGrad || i == Intrinsic::kTextureStore;
+}
+
+bool IsImageQueryIntrinsic(Intrinsic i) {
+ return i == semantic::Intrinsic::kTextureDimensions ||
+ i == Intrinsic::kTextureNumLayers ||
+ i == Intrinsic::kTextureNumLevels ||
+ i == Intrinsic::kTextureNumSamples;
+}
+
+} // namespace intrinsic
+} // namespace semantic
+} // namespace tint
diff --git a/src/semantic/intrinsic.h b/src/semantic/intrinsic.h
new file mode 100644
index 0000000..6826431
--- /dev/null
+++ b/src/semantic/intrinsic.h
@@ -0,0 +1,144 @@
+// Copyright 2020 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_INTRINSIC_H_
+#define SRC_SEMANTIC_INTRINSIC_H_
+
+#include <ostream>
+
+namespace tint {
+namespace semantic {
+
+enum class Intrinsic {
+ kNone = -1,
+
+ kAbs,
+ kAcos,
+ kAll,
+ kAny,
+ kArrayLength,
+ kAsin,
+ kAtan,
+ kAtan2,
+ kCeil,
+ kClamp,
+ kCos,
+ kCosh,
+ kCountOneBits,
+ kCross,
+ kDeterminant,
+ kDistance,
+ kDot,
+ kDpdx,
+ kDpdxCoarse,
+ kDpdxFine,
+ kDpdy,
+ kDpdyCoarse,
+ kDpdyFine,
+ kExp,
+ kExp2,
+ kFaceForward,
+ kFloor,
+ kFma,
+ kFract,
+ kFrexp,
+ kFwidth,
+ kFwidthCoarse,
+ kFwidthFine,
+ kInverseSqrt,
+ kIsFinite,
+ kIsInf,
+ kIsNan,
+ kIsNormal,
+ kLdexp,
+ kLength,
+ kLog,
+ kLog2,
+ kMax,
+ kMin,
+ kMix,
+ kModf,
+ kNormalize,
+ kPow,
+ kReflect,
+ kReverseBits,
+ kRound,
+ kSelect,
+ kSign,
+ kSin,
+ kSinh,
+ kSmoothStep,
+ kSqrt,
+ kStep,
+ kTan,
+ kTanh,
+ kTextureDimensions,
+ kTextureLoad,
+ kTextureNumLayers,
+ kTextureNumLevels,
+ kTextureNumSamples,
+ kTextureSample,
+ kTextureSampleBias,
+ kTextureSampleCompare,
+ kTextureSampleGrad,
+ kTextureSampleLevel,
+ kTextureStore,
+ kTrunc
+};
+
+/// Emits the name of the intrinsic function. The spelling,
+/// including case, matches the name in the WGSL spec.
+std::ostream& operator<<(std::ostream& out, Intrinsic i);
+
+namespace intrinsic {
+
+/// Determines if the given `i` is a coarse derivative
+/// @param i the intrinsic
+/// @returns true if the given derivative is coarse.
+bool IsCoarseDerivative(Intrinsic i);
+
+/// Determines if the given `i` is a fine derivative
+/// @param i the intrinsic
+/// @returns true if the given derivative is fine.
+bool IsFineDerivative(Intrinsic i);
+
+/// Determine if the given `i` is a derivative intrinsic
+/// @param i the intrinsic
+/// @returns true if the given `i` is a derivative intrinsic
+bool IsDerivative(Intrinsic i);
+
+/// Determines if the given `i` is a float classification intrinsic
+/// @param i the intrinsic
+/// @returns true if the given `i` is a float intrinsic
+bool IsFloatClassificationIntrinsic(Intrinsic i);
+
+/// Determines if the given `i` is a texture operation intrinsic
+/// @param i the intrinsic
+/// @returns true if the given `i` is a texture operation intrinsic
+bool IsTextureIntrinsic(Intrinsic i);
+
+/// Determines if the given `i` is a image query intrinsic
+/// @param i the intrinsic
+/// @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 semantic
+} // namespace tint
+
+#endif // SRC_SEMANTIC_INTRINSIC_H_
diff --git a/src/semantic/sem_call.cc b/src/semantic/sem_call.cc
new file mode 100644
index 0000000..fa6de41
--- /dev/null
+++ b/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
diff --git a/src/semantic/type_mappings.h b/src/semantic/type_mappings.h
index 58b91c5..95c7fff 100644
--- a/src/semantic/type_mappings.h
+++ b/src/semantic/type_mappings.h
@@ -22,6 +22,7 @@
// Forward declarations
namespace ast {
+class CallExpression;
class Expression;
class Function;
class Variable;
@@ -30,6 +31,7 @@
namespace semantic {
+class Call;
class Expression;
class Function;
class Variable;
@@ -44,6 +46,7 @@
semantic::Expression* operator()(ast::Expression*);
semantic::Function* operator()(ast::Function*);
semantic::Variable* operator()(ast::Variable*);
+ semantic::Call* operator()(ast::CallExpression*);
//! @endcond
};
diff --git a/src/type_determiner.cc b/src/type_determiner.cc
index 09cc6a6..df933c6 100644
--- a/src/type_determiner.cc
+++ b/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,56 +394,59 @@
return true;
}
-bool TypeDeterminer::DetermineCall(ast::CallExpression* expr) {
- if (!DetermineResultType(expr->func())) {
+bool TypeDeterminer::DetermineCall(ast::CallExpression* call) {
+ 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>()) {
- if (ident->IsIntrinsic()) {
- if (!DetermineIntrinsic(ident, expr)) {
- return false;
- }
- } else {
- if (current_function_) {
- caller_to_callee_[current_function_->declaration->symbol()].push_back(
- ident->symbol());
-
- auto callee_func_it = symbol_to_function_.find(ident->symbol());
- if (callee_func_it == symbol_to_function_.end()) {
- set_error(expr->source(),
- "unable to find called function: " +
- builder_->Symbols().NameFor(ident->symbol()));
- return false;
- }
- auto* callee_func = callee_func_it->second;
-
- // We inherit any referenced variables from the callee.
- for (auto* var : callee_func->referenced_module_vars) {
- set_referenced_from_function_if_needed(var, false);
- }
- }
- }
- } else {
- if (!DetermineResultType(expr->func())) {
- return false;
- }
+ auto* ident = call->func()->As<ast::IdentifierExpression>();
+ if (!ident) {
+ set_error(call->source(), "call target is not an identifier");
+ return false;
}
- if (auto* type = TypeOf(expr->func())) {
- SetType(expr, type);
+ auto name = builder_->Symbols().NameFor(ident->symbol());
+
+ auto intrinsic = MatchIntrinsic(name);
+ if (intrinsic != semantic::Intrinsic::kNone) {
+ if (!DetermineIntrinsicCall(call, intrinsic)) {
+ return false;
+ }
} else {
- auto func_sym = expr->func()->As<ast::IdentifierExpression>()->symbol();
- set_error(expr->source(),
- "v-0005: function must be declared before use: '" +
- builder_->Symbols().NameFor(func_sym) + "'");
- return false;
+ if (current_function_) {
+ caller_to_callee_[current_function_->declaration->symbol()].push_back(
+ ident->symbol());
+
+ auto callee_func_it = symbol_to_function_.find(ident->symbol());
+ if (callee_func_it == symbol_to_function_.end()) {
+ set_error(call->source(), "unable to find called function: " + name);
+ return false;
+ }
+ auto* callee_func = callee_func_it->second;
+
+ // We inherit any referenced variables from the callee.
+ for (auto* var : callee_func->referenced_module_vars) {
+ set_referenced_from_function_if_needed(var, false);
+ }
+ }
+
+ auto iter = symbol_to_function_.find(ident->symbol());
+ if (iter == symbol_to_function_.end()) {
+ set_error(call->source(),
+ "v-0005: function must be declared before use: '" + name + "'");
+ return false;
+ }
+
+ auto* function = iter->second;
+ auto* return_ty = function->declaration->return_type();
+ auto* sem = builder_->create<semantic::Call>(return_ty);
+ builder_->Sem().Add(call, sem);
}
return true;
@@ -459,7 +463,7 @@
};
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 @@
// 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},
- {ast::Intrinsic::kAcos, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kAll, IntrinsicDataType::kBool, 1, 0},
- {ast::Intrinsic::kAny, IntrinsicDataType::kBool, 1, 0},
- {ast::Intrinsic::kArrayLength, IntrinsicDataType::kUnsignedInteger, 1, 0},
- {ast::Intrinsic::kAsin, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kAtan, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kAtan2, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kCeil, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kClamp, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kCos, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kCosh, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kCountOneBits, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kCross, IntrinsicDataType::kFloat, 3, 0},
- {ast::Intrinsic::kDeterminant, IntrinsicDataType::kFloat, 1, 0},
- {ast::Intrinsic::kDistance, IntrinsicDataType::kFloat, 1, 0},
- {ast::Intrinsic::kDpdx, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kDpdxCoarse, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kDpdxFine, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kDpdy, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kDpdyCoarse, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kDpdyFine, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kDot, IntrinsicDataType::kFloat, 1, 0},
- {ast::Intrinsic::kExp, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kExp2, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kFaceForward, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kFloor, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kFwidth, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kFwidthCoarse, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kFwidthFine, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kFma, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kFract, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kFrexp, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kInverseSqrt, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kLdexp, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kLength, IntrinsicDataType::kFloat, 1, 0},
- {ast::Intrinsic::kLog, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kLog2, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kMax, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kMin, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kMix, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kModf, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kNormalize, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kPow, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kReflect, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kReverseBits, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kRound, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kSelect, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kSign, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kSin, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kSinh, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kSmoothStep, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kSqrt, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kStep, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kTan, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kTanh, IntrinsicDataType::kDependent, 0, 0},
- {ast::Intrinsic::kTrunc, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kAbs, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kAcos, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kAll, IntrinsicDataType::kBool, 1, 0},
+ {semantic::Intrinsic::kAny, IntrinsicDataType::kBool, 1, 0},
+ {semantic::Intrinsic::kArrayLength, IntrinsicDataType::kUnsignedInteger, 1,
+ 0},
+ {semantic::Intrinsic::kAsin, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kAtan, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kAtan2, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kCeil, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kClamp, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kCos, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kCosh, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kCountOneBits, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kCross, IntrinsicDataType::kFloat, 3, 0},
+ {semantic::Intrinsic::kDeterminant, IntrinsicDataType::kFloat, 1, 0},
+ {semantic::Intrinsic::kDistance, IntrinsicDataType::kFloat, 1, 0},
+ {semantic::Intrinsic::kDpdx, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kDpdxCoarse, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kDpdxFine, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kDpdy, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kDpdyCoarse, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kDpdyFine, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kDot, IntrinsicDataType::kFloat, 1, 0},
+ {semantic::Intrinsic::kExp, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kExp2, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kFaceForward, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kFloor, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kFwidth, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kFwidthCoarse, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kFwidthFine, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kFma, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kFract, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kFrexp, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kInverseSqrt, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kLdexp, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kLength, IntrinsicDataType::kFloat, 1, 0},
+ {semantic::Intrinsic::kLog, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kLog2, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kMax, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kMin, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kMix, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kModf, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kNormalize, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kPow, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kReflect, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kReverseBits, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kRound, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kSelect, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kSign, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kSin, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kSinh, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kSmoothStep, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kSqrt, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kStep, IntrinsicDataType::kDependent, 0, 0},
+ {semantic::Intrinsic::kTan, 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 @@
} // namespace
-bool TypeDeterminer::DetermineIntrinsic(ast::IdentifierExpression* ident,
- ast::CallExpression* expr) {
- if (ast::intrinsic::IsFloatClassificationIntrinsic(ident->intrinsic())) {
- if (expr->params().size() != 1) {
- set_error(expr->source(),
- "incorrect number of parameters for " +
- builder_->Symbols().NameFor(ident->symbol()));
+bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
+ semantic::Intrinsic intrinsic) {
+ auto create_sem = [&](type::Type* result) {
+ auto* sem = builder_->create<semantic::IntrinsicCall>(result, intrinsic);
+ builder_->Sem().Add(call, sem);
+ };
+
+ 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(),
- builder_->create<type::Vector>(bool_type, vec->size()));
+ create_sem(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())) {
- ast::intrinsic::TextureSignature::Parameters param;
+ if (semantic::intrinsic::IsTextureIntrinsic(intrinsic)) {
+ 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(),
- "invalid first argument for " +
- builder_->Symbols().NameFor(ident->symbol()));
+ set_error(call->source(), "invalid first argument for " + name);
return false;
}
type::Texture* texture =
@@ -568,25 +574,25 @@
bool is_array = type::IsTextureArray(texture->dim());
bool is_multisampled = texture->Is<type::MultisampledTexture>();
- switch (ident->intrinsic()) {
- case ast::Intrinsic::kTextureDimensions:
+ switch (intrinsic) {
+ 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 ast::Intrinsic::kTextureNumLevels:
- case ast::Intrinsic::kTextureNumSamples:
+ case semantic::Intrinsic::kTextureNumLayers:
+ case semantic::Intrinsic::kTextureNumLevels:
+ 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 @@
}
}
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 @@
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 @@
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 @@
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 @@
}
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 @@
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) {
- set_error(expr->source(),
- "incorrect number of parameters for " +
- builder_->Symbols().NameFor(ident->symbol()) + ", got " +
- std::to_string(expr->params().size()) + " and expected " +
+ if (call->params().size() != param.count) {
+ set_error(call->source(),
+ "incorrect number of parameters for " + name + ", got " +
+ std::to_string(call->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()) {
- case ast::Intrinsic::kTextureDimensions: {
+ switch (intrinsic) {
+ 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 @@
}
break;
}
- case ast::Intrinsic::kTextureNumLayers:
- case ast::Intrinsic::kTextureNumLevels:
- case ast::Intrinsic::kTextureNumSamples:
+ case semantic::Intrinsic::kTextureNumLayers:
+ case semantic::Intrinsic::kTextureNumLevels:
+ 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 @@
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 @@
}
}
}
- 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 " +
- builder_->Symbols().NameFor(ident->symbol());
+ error_ = "unable to find intrinsic " + name;
return false;
}
if (data->result_type == IntrinsicDataType::kDependent) {
const auto param_idx = data->param_for_result_type;
- if (expr->params().size() <= param_idx) {
- set_error(expr->source(),
+ if (call->params().size() <= param_idx) {
+ set_error(call->source(),
"missing parameter " + std::to_string(param_idx) +
- " required for type determination in builtin " +
- builder_->Symbols().NameFor(ident->symbol()));
+ " required for type determination in builtin " + name);
return false;
}
- SetType(expr->func(),
- TypeOf(expr->params()[param_idx])->UnwrapPtrIfNeeded());
+ create_sem(TypeOf(call->params()[param_idx])->UnwrapPtrIfNeeded());
} else {
// The result type is not dependent on the parameter types.
type::Type* type = nullptr;
@@ -781,15 +783,14 @@
type = builder_->create<type::Bool>();
break;
default:
- error_ = "unhandled intrinsic data type for " +
- builder_->Symbols().NameFor(ident->symbol());
+ error_ = "unhandled intrinsic data type for " + name;
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 @@
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)) {
- set_error(expr->source(),
- "v-0006: identifier must be declared before use: " +
- builder_->Symbols().NameFor(symbol));
- return false;
+ std::string name = builder_->Symbols().NameFor(symbol);
+ if (MatchIntrinsic(name) != semantic::Intrinsic::kNone) {
+ // Identifier is to an intrinsic function, which has no type (currently).
+ return true;
}
- return true;
+
+ set_error(expr->source(),
+ "v-0006: identifier must be declared before use: " + name);
+ return false;
}
-bool TypeDeterminer::SetIntrinsicIfNeeded(ast::IdentifierExpression* ident) {
- auto name = builder_->Symbols().NameFor(ident->symbol());
+semantic::Intrinsic TypeDeterminer::MatchIntrinsic(const std::string& name) {
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);
- } else {
- return false;
+ return semantic::Intrinsic::kTrunc;
}
- return true;
+ return semantic::Intrinsic::kNone;
}
bool TypeDeterminer::DetermineMemberAccessor(
diff --git a/src/type_determiner.h b/src/type_determiner.h
index 8fdf969..916a6b2 100644
--- a/src/type_determiner.h
+++ b/src/type_determiner.h
@@ -66,10 +66,10 @@
/// @returns true if the type determiner was successful
bool Determine();
- /// Sets the intrinsic data information for the identifier if needed
- /// @param ident the identifier expression
- /// @returns true if an intrinsic was set
- bool SetIntrinsicIfNeeded(ast::IdentifierExpression* ident);
+ /// @param name the function name to try and match as an intrinsic.
+ /// @return the semantic::Intrinsic for the given name. If `name` does not
+ /// match an intrinsic, returns semantic::Intrinsic::kNone
+ static semantic::Intrinsic MatchIntrinsic(const std::string& name);
private:
template <typename T>
@@ -176,8 +176,8 @@
bool DetermineCall(ast::CallExpression* expr);
bool DetermineConstructor(ast::ConstructorExpression* expr);
bool DetermineIdentifier(ast::IdentifierExpression* expr);
- bool DetermineIntrinsic(ast::IdentifierExpression* name,
- ast::CallExpression* expr);
+ bool DetermineIntrinsicCall(ast::CallExpression* call,
+ semantic::Intrinsic intrinsic);
bool DetermineMemberAccessor(ast::MemberAccessorExpression* expr);
bool DetermineUnaryOp(ast::UnaryOpExpression* expr);
diff --git a/src/type_determiner_test.cc b/src/type_determiner_test.cc
index 4878e01..a09b4f7 100644
--- a/src/type_determiner_test.cc
+++ b/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 @@
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");
- WrapInFunction(ident);
+ auto* call = Call("my_func");
+ WrapInFunction(call);
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->Is<type::F32>());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->Is<type::F32>());
}
TEST_F(TypeDeterminerTest, Expr_Identifier_Unknown) {
@@ -1521,7 +1522,7 @@
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 @@
TEST_P(IntrinsicDataTest, Lookup) {
auto param = GetParam();
- auto* ident = Expr(param.name);
- EXPECT_TRUE(td()->SetIntrinsicIfNeeded(ident));
- EXPECT_EQ(ident->intrinsic(), param.intrinsic);
- EXPECT_TRUE(ident->IsIntrinsic());
+ EXPECT_EQ(TypeDeterminer::MatchIntrinsic(param.name), param.intrinsic);
}
INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest,
IntrinsicDataTest,
testing::Values(
- IntrinsicData{"abs", ast::Intrinsic::kAbs},
- IntrinsicData{"acos", ast::Intrinsic::kAcos},
- IntrinsicData{"all", ast::Intrinsic::kAll},
- IntrinsicData{"any", ast::Intrinsic::kAny},
- IntrinsicData{"arrayLength", ast::Intrinsic::kArrayLength},
- IntrinsicData{"asin", ast::Intrinsic::kAsin},
- IntrinsicData{"atan", ast::Intrinsic::kAtan},
- IntrinsicData{"atan2", ast::Intrinsic::kAtan2},
- IntrinsicData{"ceil", ast::Intrinsic::kCeil},
- IntrinsicData{"clamp", ast::Intrinsic::kClamp},
- IntrinsicData{"cos", ast::Intrinsic::kCos},
- IntrinsicData{"cosh", ast::Intrinsic::kCosh},
- IntrinsicData{"countOneBits", ast::Intrinsic::kCountOneBits},
- IntrinsicData{"cross", ast::Intrinsic::kCross},
- IntrinsicData{"determinant", ast::Intrinsic::kDeterminant},
- IntrinsicData{"distance", ast::Intrinsic::kDistance},
- IntrinsicData{"dot", ast::Intrinsic::kDot},
- IntrinsicData{"dpdx", ast::Intrinsic::kDpdx},
- IntrinsicData{"dpdxCoarse", ast::Intrinsic::kDpdxCoarse},
- IntrinsicData{"dpdxFine", ast::Intrinsic::kDpdxFine},
- IntrinsicData{"dpdy", ast::Intrinsic::kDpdy},
- IntrinsicData{"dpdyCoarse", ast::Intrinsic::kDpdyCoarse},
- IntrinsicData{"dpdyFine", ast::Intrinsic::kDpdyFine},
- IntrinsicData{"exp", ast::Intrinsic::kExp},
- IntrinsicData{"exp2", ast::Intrinsic::kExp2},
- IntrinsicData{"faceForward", ast::Intrinsic::kFaceForward},
- IntrinsicData{"floor", ast::Intrinsic::kFloor},
- IntrinsicData{"fma", ast::Intrinsic::kFma},
- IntrinsicData{"fract", ast::Intrinsic::kFract},
- IntrinsicData{"frexp", ast::Intrinsic::kFrexp},
- IntrinsicData{"fwidth", ast::Intrinsic::kFwidth},
- IntrinsicData{"fwidthCoarse", ast::Intrinsic::kFwidthCoarse},
- IntrinsicData{"fwidthFine", ast::Intrinsic::kFwidthFine},
- IntrinsicData{"inverseSqrt", ast::Intrinsic::kInverseSqrt},
- IntrinsicData{"isFinite", ast::Intrinsic::kIsFinite},
- IntrinsicData{"isInf", ast::Intrinsic::kIsInf},
- IntrinsicData{"isNan", ast::Intrinsic::kIsNan},
- IntrinsicData{"isNormal", ast::Intrinsic::kIsNormal},
- IntrinsicData{"ldexp", ast::Intrinsic::kLdexp},
- IntrinsicData{"length", ast::Intrinsic::kLength},
- IntrinsicData{"log", ast::Intrinsic::kLog},
- IntrinsicData{"log2", ast::Intrinsic::kLog2},
- IntrinsicData{"max", ast::Intrinsic::kMax},
- IntrinsicData{"min", ast::Intrinsic::kMin},
- IntrinsicData{"mix", ast::Intrinsic::kMix},
- IntrinsicData{"modf", ast::Intrinsic::kModf},
- IntrinsicData{"normalize", ast::Intrinsic::kNormalize},
- IntrinsicData{"pow", ast::Intrinsic::kPow},
- IntrinsicData{"reflect", ast::Intrinsic::kReflect},
- IntrinsicData{"reverseBits", ast::Intrinsic::kReverseBits},
- IntrinsicData{"round", ast::Intrinsic::kRound},
- IntrinsicData{"select", ast::Intrinsic::kSelect},
- IntrinsicData{"sign", ast::Intrinsic::kSign},
- IntrinsicData{"sin", ast::Intrinsic::kSin},
- IntrinsicData{"sinh", ast::Intrinsic::kSinh},
- IntrinsicData{"smoothStep", ast::Intrinsic::kSmoothStep},
- IntrinsicData{"sqrt", ast::Intrinsic::kSqrt},
- IntrinsicData{"step", ast::Intrinsic::kStep},
- IntrinsicData{"tan", ast::Intrinsic::kTan},
- IntrinsicData{"tanh", ast::Intrinsic::kTanh},
- IntrinsicData{"textureDimensions", ast::Intrinsic::kTextureDimensions},
- IntrinsicData{"textureLoad", ast::Intrinsic::kTextureLoad},
- IntrinsicData{"textureNumLayers", ast::Intrinsic::kTextureNumLayers},
- IntrinsicData{"textureNumLevels", ast::Intrinsic::kTextureNumLevels},
- IntrinsicData{"textureNumSamples", ast::Intrinsic::kTextureNumSamples},
- IntrinsicData{"textureSample", ast::Intrinsic::kTextureSample},
- IntrinsicData{"textureSampleBias", ast::Intrinsic::kTextureSampleBias},
+ IntrinsicData{"abs", semantic::Intrinsic::kAbs},
+ IntrinsicData{"acos", semantic::Intrinsic::kAcos},
+ IntrinsicData{"all", semantic::Intrinsic::kAll},
+ IntrinsicData{"any", semantic::Intrinsic::kAny},
+ IntrinsicData{"arrayLength", semantic::Intrinsic::kArrayLength},
+ IntrinsicData{"asin", semantic::Intrinsic::kAsin},
+ IntrinsicData{"atan", semantic::Intrinsic::kAtan},
+ IntrinsicData{"atan2", semantic::Intrinsic::kAtan2},
+ IntrinsicData{"ceil", semantic::Intrinsic::kCeil},
+ IntrinsicData{"clamp", semantic::Intrinsic::kClamp},
+ IntrinsicData{"cos", semantic::Intrinsic::kCos},
+ IntrinsicData{"cosh", semantic::Intrinsic::kCosh},
+ IntrinsicData{"countOneBits", semantic::Intrinsic::kCountOneBits},
+ IntrinsicData{"cross", semantic::Intrinsic::kCross},
+ IntrinsicData{"determinant", semantic::Intrinsic::kDeterminant},
+ IntrinsicData{"distance", semantic::Intrinsic::kDistance},
+ IntrinsicData{"dot", semantic::Intrinsic::kDot},
+ IntrinsicData{"dpdx", semantic::Intrinsic::kDpdx},
+ IntrinsicData{"dpdxCoarse", semantic::Intrinsic::kDpdxCoarse},
+ IntrinsicData{"dpdxFine", semantic::Intrinsic::kDpdxFine},
+ IntrinsicData{"dpdy", semantic::Intrinsic::kDpdy},
+ IntrinsicData{"dpdyCoarse", semantic::Intrinsic::kDpdyCoarse},
+ IntrinsicData{"dpdyFine", semantic::Intrinsic::kDpdyFine},
+ IntrinsicData{"exp", semantic::Intrinsic::kExp},
+ IntrinsicData{"exp2", semantic::Intrinsic::kExp2},
+ IntrinsicData{"faceForward", semantic::Intrinsic::kFaceForward},
+ IntrinsicData{"floor", semantic::Intrinsic::kFloor},
+ IntrinsicData{"fma", semantic::Intrinsic::kFma},
+ IntrinsicData{"fract", semantic::Intrinsic::kFract},
+ IntrinsicData{"frexp", semantic::Intrinsic::kFrexp},
+ IntrinsicData{"fwidth", semantic::Intrinsic::kFwidth},
+ IntrinsicData{"fwidthCoarse", semantic::Intrinsic::kFwidthCoarse},
+ IntrinsicData{"fwidthFine", semantic::Intrinsic::kFwidthFine},
+ IntrinsicData{"inverseSqrt", semantic::Intrinsic::kInverseSqrt},
+ IntrinsicData{"isFinite", semantic::Intrinsic::kIsFinite},
+ IntrinsicData{"isInf", semantic::Intrinsic::kIsInf},
+ IntrinsicData{"isNan", semantic::Intrinsic::kIsNan},
+ IntrinsicData{"isNormal", semantic::Intrinsic::kIsNormal},
+ IntrinsicData{"ldexp", semantic::Intrinsic::kLdexp},
+ IntrinsicData{"length", semantic::Intrinsic::kLength},
+ IntrinsicData{"log", semantic::Intrinsic::kLog},
+ IntrinsicData{"log2", semantic::Intrinsic::kLog2},
+ IntrinsicData{"max", semantic::Intrinsic::kMax},
+ IntrinsicData{"min", semantic::Intrinsic::kMin},
+ IntrinsicData{"mix", semantic::Intrinsic::kMix},
+ IntrinsicData{"modf", semantic::Intrinsic::kModf},
+ IntrinsicData{"normalize", semantic::Intrinsic::kNormalize},
+ IntrinsicData{"pow", semantic::Intrinsic::kPow},
+ IntrinsicData{"reflect", semantic::Intrinsic::kReflect},
+ IntrinsicData{"reverseBits", semantic::Intrinsic::kReverseBits},
+ IntrinsicData{"round", semantic::Intrinsic::kRound},
+ IntrinsicData{"select", semantic::Intrinsic::kSelect},
+ IntrinsicData{"sign", semantic::Intrinsic::kSign},
+ IntrinsicData{"sin", semantic::Intrinsic::kSin},
+ IntrinsicData{"sinh", semantic::Intrinsic::kSinh},
+ IntrinsicData{"smoothStep", semantic::Intrinsic::kSmoothStep},
+ IntrinsicData{"sqrt", semantic::Intrinsic::kSqrt},
+ IntrinsicData{"step", semantic::Intrinsic::kStep},
+ IntrinsicData{"tan", semantic::Intrinsic::kTan},
+ IntrinsicData{"tanh", semantic::Intrinsic::kTanh},
+ IntrinsicData{"textureDimensions",
+ semantic::Intrinsic::kTextureDimensions},
+ IntrinsicData{"textureLoad", semantic::Intrinsic::kTextureLoad},
+ IntrinsicData{"textureNumLayers",
+ semantic::Intrinsic::kTextureNumLayers},
+ IntrinsicData{"textureNumLevels",
+ semantic::Intrinsic::kTextureNumLevels},
+ IntrinsicData{"textureNumSamples",
+ semantic::Intrinsic::kTextureNumSamples},
+ IntrinsicData{"textureSample", semantic::Intrinsic::kTextureSample},
+ IntrinsicData{"textureSampleBias",
+ semantic::Intrinsic::kTextureSampleBias},
IntrinsicData{"textureSampleCompare",
- ast::Intrinsic::kTextureSampleCompare},
- IntrinsicData{"textureSampleGrad", ast::Intrinsic::kTextureSampleGrad},
+ semantic::Intrinsic::kTextureSampleCompare},
+ IntrinsicData{"textureSampleGrad",
+ semantic::Intrinsic::kTextureSampleGrad},
IntrinsicData{"textureSampleLevel",
- ast::Intrinsic::kTextureSampleLevel},
- IntrinsicData{"trunc", ast::Intrinsic::kTrunc}));
+ semantic::Intrinsic::kTextureSampleLevel},
+ IntrinsicData{"trunc", semantic::Intrinsic::kTrunc}));
-TEST_F(TypeDeterminerTest, IntrinsicNotSetIfNotMatched) {
- auto* ident = Expr("not_intrinsic");
- EXPECT_FALSE(td()->SetIntrinsicIfNeeded(ident));
- EXPECT_EQ(ident->intrinsic(), ast::Intrinsic::kNone);
- EXPECT_FALSE(ident->IsIntrinsic());
+TEST_F(TypeDeterminerTest, MatchIntrinsicNoMatch) {
+ EXPECT_EQ(TypeDeterminer::MatchIntrinsic("not_intrinsic"),
+ semantic::Intrinsic::kNone);
}
using ImportData_SingleParamTest = TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1631,8 +1633,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_scalar());
}
TEST_P(ImportData_SingleParamTest, Vector) {
@@ -1644,9 +1646,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_SingleParamTest, Error_NoParams) {
@@ -1665,28 +1667,29 @@
INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest,
ImportData_SingleParamTest,
- testing::Values(IntrinsicData{"acos", ast::Intrinsic::kAcos},
- IntrinsicData{"asin", ast::Intrinsic::kAsin},
- IntrinsicData{"atan", ast::Intrinsic::kAtan},
- IntrinsicData{"ceil", ast::Intrinsic::kCeil},
- IntrinsicData{"cos", ast::Intrinsic::kCos},
- IntrinsicData{"cosh", ast::Intrinsic::kCosh},
- IntrinsicData{"exp", ast::Intrinsic::kExp},
- IntrinsicData{"exp2", ast::Intrinsic::kExp2},
- IntrinsicData{"floor", ast::Intrinsic::kFloor},
- IntrinsicData{"fract", ast::Intrinsic::kFract},
- IntrinsicData{"inverseSqrt", ast::Intrinsic::kInverseSqrt},
- IntrinsicData{"log", ast::Intrinsic::kLog},
- IntrinsicData{"log2", ast::Intrinsic::kLog2},
- IntrinsicData{"normalize", ast::Intrinsic::kNormalize},
- IntrinsicData{"round", ast::Intrinsic::kRound},
- IntrinsicData{"sign", ast::Intrinsic::kSign},
- IntrinsicData{"sin", ast::Intrinsic::kSin},
- IntrinsicData{"sinh", ast::Intrinsic::kSinh},
- IntrinsicData{"sqrt", ast::Intrinsic::kSqrt},
- IntrinsicData{"tan", ast::Intrinsic::kTan},
- IntrinsicData{"tanh", ast::Intrinsic::kTanh},
- IntrinsicData{"trunc", ast::Intrinsic::kTrunc}));
+ testing::Values(IntrinsicData{"acos", semantic::Intrinsic::kAcos},
+ IntrinsicData{"asin", semantic::Intrinsic::kAsin},
+ IntrinsicData{"atan", semantic::Intrinsic::kAtan},
+ IntrinsicData{"ceil", semantic::Intrinsic::kCeil},
+ IntrinsicData{"cos", semantic::Intrinsic::kCos},
+ IntrinsicData{"cosh", semantic::Intrinsic::kCosh},
+ IntrinsicData{"exp", semantic::Intrinsic::kExp},
+ IntrinsicData{"exp2", semantic::Intrinsic::kExp2},
+ IntrinsicData{"floor", semantic::Intrinsic::kFloor},
+ IntrinsicData{"fract", semantic::Intrinsic::kFract},
+ IntrinsicData{"inverseSqrt",
+ semantic::Intrinsic::kInverseSqrt},
+ IntrinsicData{"log", semantic::Intrinsic::kLog},
+ IntrinsicData{"log2", semantic::Intrinsic::kLog2},
+ IntrinsicData{"normalize", semantic::Intrinsic::kNormalize},
+ IntrinsicData{"round", semantic::Intrinsic::kRound},
+ IntrinsicData{"sign", semantic::Intrinsic::kSign},
+ IntrinsicData{"sin", semantic::Intrinsic::kSin},
+ IntrinsicData{"sinh", semantic::Intrinsic::kSinh},
+ IntrinsicData{"sqrt", semantic::Intrinsic::kSqrt},
+ IntrinsicData{"tan", semantic::Intrinsic::kTan},
+ IntrinsicData{"tanh", semantic::Intrinsic::kTanh},
+ IntrinsicData{"trunc", semantic::Intrinsic::kTrunc}));
using ImportData_SingleParam_FloatOrInt_Test =
TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1699,8 +1702,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_scalar());
}
TEST_P(ImportData_SingleParam_FloatOrInt_Test, Float_Vector) {
@@ -1712,9 +1715,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_SingleParam_FloatOrInt_Test, Sint_Scalar) {
@@ -1726,8 +1729,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->Is<type::I32>());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->Is<type::I32>());
}
TEST_P(ImportData_SingleParam_FloatOrInt_Test, Sint_Vector) {
@@ -1747,9 +1750,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_signed_integer_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_signed_integer_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_SingleParam_FloatOrInt_Test, Uint_Scalar) {
@@ -1764,8 +1767,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->Is<type::U32>());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->Is<type::U32>());
}
TEST_P(ImportData_SingleParam_FloatOrInt_Test, Uint_Vector) {
@@ -1777,9 +1780,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_unsigned_integer_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_unsigned_integer_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_SingleParam_FloatOrInt_Test, Error_NoParams) {
@@ -1797,8 +1800,8 @@
INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest,
ImportData_SingleParam_FloatOrInt_Test,
- testing::Values(IntrinsicData{"abs",
- ast::Intrinsic::kAbs}));
+ testing::Values(IntrinsicData{
+ "abs", semantic::Intrinsic::kAbs}));
TEST_F(TypeDeterminerTest, ImportData_Length_Scalar) {
auto* ident = Expr("length");
@@ -1808,8 +1811,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_scalar());
}
TEST_F(TypeDeterminerTest, ImportData_Length_FloatVector) {
@@ -1823,8 +1826,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_scalar());
}
using ImportData_TwoParamTest = TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1837,8 +1840,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_scalar());
}
TEST_P(ImportData_TwoParamTest, Vector) {
@@ -1851,9 +1854,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_TwoParamTest, Error_NoParams) {
auto param = GetParam();
@@ -1870,10 +1873,10 @@
INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest,
ImportData_TwoParamTest,
- testing::Values(IntrinsicData{"atan2", ast::Intrinsic::kAtan2},
- IntrinsicData{"pow", ast::Intrinsic::kPow},
- IntrinsicData{"step", ast::Intrinsic::kStep},
- IntrinsicData{"reflect", ast::Intrinsic::kReflect}));
+ testing::Values(IntrinsicData{"atan2", semantic::Intrinsic::kAtan2},
+ IntrinsicData{"pow", semantic::Intrinsic::kPow},
+ IntrinsicData{"step", semantic::Intrinsic::kStep},
+ IntrinsicData{"reflect", semantic::Intrinsic::kReflect}));
TEST_F(TypeDeterminerTest, ImportData_Distance_Scalar) {
auto* ident = Expr("distance");
@@ -1883,8 +1886,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_scalar());
}
TEST_F(TypeDeterminerTest, ImportData_Distance_Vector) {
@@ -1896,8 +1899,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->Is<type::F32>());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->Is<type::F32>());
}
TEST_F(TypeDeterminerTest, ImportData_Cross) {
@@ -1909,9 +1912,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_F(TypeDeterminerTest, ImportData_Cross_AutoType) {
@@ -1922,9 +1925,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
using ImportData_ThreeParamTest = TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1937,8 +1940,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_scalar());
}
TEST_P(ImportData_ThreeParamTest, Vector) {
@@ -1951,9 +1954,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_ThreeParamTest, Error_NoParams) {
auto param = GetParam();
@@ -1971,11 +1974,11 @@
INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest,
ImportData_ThreeParamTest,
- testing::Values(IntrinsicData{"mix", ast::Intrinsic::kMix},
- IntrinsicData{"smoothStep", ast::Intrinsic::kSmoothStep},
- IntrinsicData{"fma", ast::Intrinsic::kFma},
- IntrinsicData{"faceForward",
- ast::Intrinsic::kFaceForward}));
+ testing::Values(
+ IntrinsicData{"mix", semantic::Intrinsic::kMix},
+ IntrinsicData{"smoothStep", semantic::Intrinsic::kSmoothStep},
+ IntrinsicData{"fma", semantic::Intrinsic::kFma},
+ IntrinsicData{"faceForward", semantic::Intrinsic::kFaceForward}));
using ImportData_ThreeParam_FloatOrInt_Test =
TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1988,8 +1991,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_scalar());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_scalar());
}
TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Float_Vector) {
@@ -2002,9 +2005,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Sint_Scalar) {
@@ -2016,8 +2019,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->Is<type::I32>());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->Is<type::I32>());
}
TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Sint_Vector) {
@@ -2030,9 +2033,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_signed_integer_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_signed_integer_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Uint_Scalar) {
@@ -2044,8 +2047,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->Is<type::U32>());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->Is<type::U32>());
}
TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Uint_Vector) {
@@ -2058,9 +2061,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_unsigned_integer_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_unsigned_integer_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Error_NoParams) {
@@ -2079,7 +2082,7 @@
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 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_integer_scalar());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_integer_scalar());
}
TEST_P(ImportData_Int_SingleParamTest, Vector) {
@@ -2105,9 +2108,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_signed_integer_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_signed_integer_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_Int_SingleParamTest, Error_NoParams) {
@@ -2127,8 +2130,8 @@
TypeDeterminerTest,
ImportData_Int_SingleParamTest,
testing::Values(
- IntrinsicData{"countOneBits", ast::Intrinsic::kCountOneBits},
- IntrinsicData{"reverseBits", ast::Intrinsic::kReverseBits}));
+ IntrinsicData{"countOneBits", semantic::Intrinsic::kCountOneBits},
+ IntrinsicData{"reverseBits", semantic::Intrinsic::kReverseBits}));
using ImportData_FloatOrInt_TwoParamTest =
TypeDeterminerTestWithParam<IntrinsicData>;
@@ -2141,8 +2144,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->Is<type::I32>());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->Is<type::I32>());
}
TEST_P(ImportData_FloatOrInt_TwoParamTest, Scalar_Unsigned) {
@@ -2154,8 +2157,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->Is<type::U32>());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->Is<type::U32>());
}
TEST_P(ImportData_FloatOrInt_TwoParamTest, Scalar_Float) {
@@ -2167,8 +2170,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->Is<type::F32>());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->Is<type::F32>());
}
TEST_P(ImportData_FloatOrInt_TwoParamTest, Vector_Signed) {
@@ -2180,9 +2183,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_signed_integer_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_signed_integer_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_FloatOrInt_TwoParamTest, Vector_Unsigned) {
@@ -2194,9 +2197,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_unsigned_integer_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_unsigned_integer_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_FloatOrInt_TwoParamTest, Vector_Float) {
@@ -2208,9 +2211,9 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->is_float_vector());
- EXPECT_EQ(TypeOf(ident)->As<type::Vector>()->size(), 3u);
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->is_float_vector());
+ EXPECT_EQ(TypeOf(call)->As<type::Vector>()->size(), 3u);
}
TEST_P(ImportData_FloatOrInt_TwoParamTest, Error_NoParams) {
@@ -2229,8 +2232,8 @@
INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest,
ImportData_FloatOrInt_TwoParamTest,
- testing::Values(IntrinsicData{"min", ast::Intrinsic::kMin},
- IntrinsicData{"max", ast::Intrinsic::kMax}));
+ testing::Values(IntrinsicData{"min", semantic::Intrinsic::kMin},
+ IntrinsicData{"max", semantic::Intrinsic::kMax}));
TEST_F(TypeDeterminerTest, ImportData_GLSL_Determinant) {
Global("var", ast::StorageClass::kFunction, ty.mat3x3<f32>());
@@ -2242,8 +2245,8 @@
EXPECT_TRUE(td()->Determine()) << td()->error();
- ASSERT_NE(TypeOf(ident), nullptr);
- EXPECT_TRUE(TypeOf(ident)->Is<type::F32>());
+ ASSERT_NE(TypeOf(call), nullptr);
+ EXPECT_TRUE(TypeOf(call)->Is<type::F32>());
}
using ImportData_Matrix_OneParam_Test =
@@ -2263,7 +2266,8 @@
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 @@
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;
- auto maybe_add_param = [¶ms](size_t idx, const char* name) {
- if (idx != ast::intrinsic::TextureSignature::Parameters::kNotUsed) {
- params.emplace_back(Parameter{idx, name});
+ std::vector<Parameter> list;
+ auto maybe_add_param = [&list](size_t idx, const char* name) {
+ if (idx !=
+ semantic::TextureIntrinsicCall::Parameters::Parameters::kNotUsed) {
+ list.emplace_back(Parameter{idx, name});
}
};
- maybe_add_param(sig->params.idx.array_index, "array_index");
- maybe_add_param(sig->params.idx.bias, "bias");
- maybe_add_param(sig->params.idx.coords, "coords");
- maybe_add_param(sig->params.idx.depth_ref, "depth_ref");
- maybe_add_param(sig->params.idx.ddx, "ddx");
- maybe_add_param(sig->params.idx.ddy, "ddy");
- maybe_add_param(sig->params.idx.level, "level");
- maybe_add_param(sig->params.idx.offset, "offset");
- maybe_add_param(sig->params.idx.sampler, "sampler");
- maybe_add_param(sig->params.idx.sample_index, "sample_index");
- maybe_add_param(sig->params.idx.texture, "texture");
- maybe_add_param(sig->params.idx.value, "value");
+ maybe_add_param(params.idx.array_index, "array_index");
+ maybe_add_param(params.idx.bias, "bias");
+ maybe_add_param(params.idx.coords, "coords");
+ maybe_add_param(params.idx.depth_ref, "depth_ref");
+ maybe_add_param(params.idx.ddx, "ddx");
+ maybe_add_param(params.idx.ddy, "ddy");
+ maybe_add_param(params.idx.level, "level");
+ maybe_add_param(params.idx.offset, "offset");
+ maybe_add_param(params.idx.sampler, "sampler");
+ maybe_add_param(params.idx.sample_index, "sample_index");
+ maybe_add_param(params.idx.texture, "texture");
+ 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 @@
}
}
- auto* sig = static_cast<const ast::intrinsic::TextureSignature*>(
- ident->intrinsic_signature());
- ASSERT_NE(sig, nullptr);
+ auto* sem = Sem().Get(call);
+ ASSERT_NE(sem, 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);
}
diff --git a/src/validator/validator_impl.cc b/src/validator/validator_impl.cc
index 265ad84..e1c999f 100644
--- a/src/validator/validator_impl.cc
+++ b/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 @@
};
struct IntrinsicData {
- ast::Intrinsic intrinsic;
+ semantic::Intrinsic intrinsic;
uint32_t param_count;
IntrinsicDataType data_type;
uint32_t vector_size;
@@ -72,102 +73,112 @@
// 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},
- {ast::Intrinsic::kAny, 1, IntrinsicDataType::kBoolVector, 0, false},
- {ast::Intrinsic::kArrayLength, 1, IntrinsicDataType::kMixed, 0, false},
- {ast::Intrinsic::kAsin, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {semantic::Intrinsic::kAtan, 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},
- {ast::Intrinsic::kCosh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {ast::Intrinsic::kCountOneBits, 1, IntrinsicDataType::kIntScalarOrVector, 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,
+ {semantic::Intrinsic::kCountOneBits, 1,
+ IntrinsicDataType::kIntScalarOrVector, 0, true},
+ {semantic::Intrinsic::kCross, 2, IntrinsicDataType::kFloatVector, 3, true},
+ {semantic::Intrinsic::kDeterminant, 1, IntrinsicDataType::kMatrix, 0,
false},
- {ast::Intrinsic::kDot, 2, IntrinsicDataType::kFloatVector, 0, false},
- {ast::Intrinsic::kDpdx, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {semantic::Intrinsic::kDistance, 2, IntrinsicDataType::kFloatScalarOrVector,
+ 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,
- true},
- {ast::Intrinsic::kDpdxFine, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- 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,
+ {semantic::Intrinsic::kDpdxCoarse, 1,
+ IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {semantic::Intrinsic::kDpdxFine, 1, 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},
- {ast::Intrinsic::kFract, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {ast::Intrinsic::kFrexp, 2, IntrinsicDataType::kMixed, 0, false},
- {ast::Intrinsic::kFwidth, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {ast::Intrinsic::kFwidthCoarse, 1, IntrinsicDataType::kFloatScalarOrVector,
+ {semantic::Intrinsic::kDpdyCoarse, 1,
+ IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {semantic::Intrinsic::kDpdyFine, 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,
- false},
- {ast::Intrinsic::kLog, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
- {ast::Intrinsic::kLog2, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {semantic::Intrinsic::kLength, 1, IntrinsicDataType::kFloatScalarOrVector,
+ 0, false},
+ {semantic::Intrinsic::kLog, 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},
- {ast::Intrinsic::kModf, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {semantic::Intrinsic::kModf, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
true},
- {ast::Intrinsic::kNormalize, 1, IntrinsicDataType::kFloatVector, 0, true},
- {ast::Intrinsic::kPow, 2, IntrinsicDataType::kFloatScalarOrVector, 0, true},
- {ast::Intrinsic::kReflect, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {semantic::Intrinsic::kNormalize, 1, IntrinsicDataType::kFloatVector, 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},
- {ast::Intrinsic::kSign, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {semantic::Intrinsic::kSelect, 3, IntrinsicDataType::kMixed, 0, false},
+ {semantic::Intrinsic::kSign, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true},
- {ast::Intrinsic::kSin, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
- {ast::Intrinsic::kSinh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {semantic::Intrinsic::kSin, 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},
- {ast::Intrinsic::kTanh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {semantic::Intrinsic::kTan, 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,248 +657,245 @@
return false;
}
- if (auto* ident = expr->func()->As<ast::IdentifierExpression>()) {
- auto symbol = ident->symbol();
- if (ident->IsIntrinsic()) {
- const IntrinsicData* data = nullptr;
- for (uint32_t i = 0; i < kIntrinsicDataCount; ++i) {
- if (ident->intrinsic() == kIntrinsicData[i].intrinsic) {
- data = &kIntrinsicData[i];
- break;
- }
+ auto* call_sem = program_->Sem().Get(expr);
+ if (call_sem == nullptr) {
+ 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 (intrinsic_sem->intrinsic() == kIntrinsicData[i].intrinsic) {
+ data = &kIntrinsicData[i];
+ break;
+ }
+ }
+
+ if (data != nullptr) {
+ 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 +
+ " expected " + std::to_string(data->param_count) +
+ " got " + std::to_string(expr->params().size()));
+ return false;
}
- if (data != nullptr) {
- const auto builtin = program_->Symbols().NameFor(symbol);
- if (expr->params().size() != data->param_count) {
- add_error(expr->source(),
- "incorrect number of parameters for " + builtin +
- " expected " + std::to_string(data->param_count) +
- " got " + std::to_string(expr->params().size()));
+ if (data->all_types_match) {
+ // Check that the type is an acceptable one.
+ if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
+ data->data_type, data->vector_size, this)) {
return false;
}
- if (data->all_types_match) {
- // Check that the type is an acceptable one.
- if (!IsValidType(program_->TypeOf(expr->func()), expr->source(),
- builtin, data->data_type, data->vector_size, this)) {
+ // Check that all params match the result type.
+ for (uint32_t i = 0; i < data->param_count; ++i) {
+ if (program_->TypeOf(expr)->UnwrapPtrIfNeeded() !=
+ program_->TypeOf(expr->params()[i])->UnwrapPtrIfNeeded()) {
+ add_error(expr->params()[i]->source(),
+ "expected parameter " + std::to_string(i) +
+ "'s unwrapped type to match result type for " +
+ builtin);
+ return false;
+ }
+ }
+ } else {
+ if (data->data_type != IntrinsicDataType::kMixed) {
+ auto* p0 = expr->params()[0];
+ if (!IsValidType(program_->TypeOf(p0), p0->source(), builtin,
+ 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() !=
+ // Check that parameters are valid types.
+ for (uint32_t i = 1; i < expr->params().size(); ++i) {
+ if (program_->TypeOf(p0)->UnwrapPtrIfNeeded() !=
program_->TypeOf(expr->params()[i])->UnwrapPtrIfNeeded()) {
- add_error(expr->params()[i]->source(),
- "expected parameter " + std::to_string(i) +
- "'s unwrapped type to match result type for " +
- builtin);
+ add_error(expr->source(),
+ "parameter " + std::to_string(i) +
+ "'s unwrapped type must match parameter 0's type");
return false;
}
}
} else {
- if (data->data_type != IntrinsicDataType::kMixed) {
+ // Special cases.
+ if (data->intrinsic == semantic::Intrinsic::kFrexp) {
auto* p0 = expr->params()[0];
- if (!IsValidType(program_->TypeOf(p0), p0->source(), builtin,
- data->data_type, data->vector_size, this)) {
+ auto* p1 = expr->params()[1];
+ auto* t0 = program_->TypeOf(p0)->UnwrapPtrIfNeeded();
+ auto* t1 = program_->TypeOf(p1)->UnwrapPtrIfNeeded();
+ if (!IsValidType(t0, p0->source(), builtin,
+ IntrinsicDataType::kFloatScalarOrVector, 0,
+ this)) {
+ return false;
+ }
+ if (!IsValidType(t1, p1->source(), builtin,
+ IntrinsicDataType::kIntScalarOrVector, 0, this)) {
return false;
}
- // Check that parameters are valid types.
- for (uint32_t i = 1; i < expr->params().size(); ++i) {
- if (program_->TypeOf(p0)->UnwrapPtrIfNeeded() !=
- program_->TypeOf(expr->params()[i])->UnwrapPtrIfNeeded()) {
+ if (t0->is_scalar()) {
+ if (!t1->is_scalar()) {
add_error(
expr->source(),
- "parameter " + std::to_string(i) +
- "'s unwrapped type must match parameter 0's type");
+ "incorrect types for " + builtin +
+ ". Parameters must be matched scalars or vectors");
return false;
}
- }
- } else {
- // Special cases.
- if (data->intrinsic == ast::Intrinsic::kFrexp) {
- auto* p0 = expr->params()[0];
- auto* p1 = expr->params()[1];
- auto* t0 = program_->TypeOf(p0)->UnwrapPtrIfNeeded();
- auto* t1 = program_->TypeOf(p1)->UnwrapPtrIfNeeded();
- if (!IsValidType(t0, p0->source(), builtin,
- IntrinsicDataType::kFloatScalarOrVector, 0,
- this)) {
- return false;
- }
- if (!IsValidType(t1, p1->source(), builtin,
- IntrinsicDataType::kIntScalarOrVector, 0,
- this)) {
- return false;
- }
-
- if (t0->is_scalar()) {
- if (!t1->is_scalar()) {
- add_error(
- expr->source(),
- "incorrect types for " + builtin +
- ". Parameters must be matched scalars or vectors");
- return false;
- }
- } else {
- if (t1->is_integer_scalar()) {
- add_error(
- expr->source(),
- "incorrect types for " + builtin +
- ". Parameters must be matched scalars or vectors");
- return false;
- }
- const auto* v0 = t0->As<type::Vector>();
- const auto* v1 = t1->As<type::Vector>();
- if (v0->size() != v1->size()) {
- add_error(expr->source(),
- "incorrect types for " + builtin +
- ". Parameter vector sizes must match");
- return false;
- }
- }
- }
-
- if (data->intrinsic == ast::Intrinsic::kSelect) {
- auto* type = program_->TypeOf(expr->func());
- auto* t0 =
- program_->TypeOf(expr->params()[0])->UnwrapPtrIfNeeded();
- auto* t1 =
- 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 +
- ". Requires bool, int or float scalar or vector");
- return false;
- }
-
- if (type != t0 || type != t1) {
- add_error(expr->source(),
- "incorrect type for " + builtin +
- ". Value parameter types must match result type");
- return false;
- }
-
- if (!t2->is_bool_scalar_or_vector()) {
- add_error(
- expr->params()[2]->source(),
- "incorrect type for " + builtin +
- ". Selector must be a bool scalar or vector value");
- return false;
- }
-
- if (type->Is<type::Vector>()) {
- auto size = type->As<type::Vector>()->size();
- if (t2->is_scalar() || size != t2->As<type::Vector>()->size()) {
- add_error(expr->params()[2]->source(),
- "incorrect type for " + builtin +
- ". Selector must be a vector with the same "
- "number of elements as the result type");
- return false;
- }
- } else {
- if (!t2->is_scalar()) {
- add_error(expr->params()[2]->source(),
- "incorrect type for " + builtin +
- ". Selector must be a bool scalar to match "
- "scalar result type");
- return false;
- }
- }
- }
-
- if (data->intrinsic == ast::Intrinsic::kArrayLength) {
- if (!program_->TypeOf(expr->func())
- ->UnwrapPtrIfNeeded()
- ->Is<type::U32>()) {
+ } else {
+ if (t1->is_integer_scalar()) {
add_error(
expr->source(),
- "incorrect type for " + builtin +
- ". Result type must be an unsigned int scalar value");
+ "incorrect types for " + builtin +
+ ". Parameters must be matched scalars or vectors");
return false;
}
+ const auto* v0 = t0->As<type::Vector>();
+ const auto* v1 = t1->As<type::Vector>();
+ if (v0->size() != v1->size()) {
+ add_error(expr->source(),
+ "incorrect types for " + builtin +
+ ". Parameter vector sizes must match");
+ return false;
+ }
+ }
+ }
- auto* p0 =
- program_->TypeOf(expr->params()[0])->UnwrapPtrIfNeeded();
- if (!p0->Is<type::Array>() ||
- !p0->As<type::Array>()->IsRuntimeArray()) {
- add_error(expr->params()[0]->source(),
+ if (data->intrinsic == semantic::Intrinsic::kSelect) {
+ auto* type = program_->TypeOf(expr);
+ auto* t0 = program_->TypeOf(expr->params()[0])->UnwrapPtrIfNeeded();
+ auto* t1 = 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 +
+ ". Requires bool, int or float scalar or vector");
+ return false;
+ }
+
+ if (type != t0 || type != t1) {
+ add_error(expr->source(),
+ "incorrect type for " + builtin +
+ ". Value parameter types must match result type");
+ return false;
+ }
+
+ if (!t2->is_bool_scalar_or_vector()) {
+ add_error(expr->params()[2]->source(),
+ "incorrect type for " + builtin +
+ ". Selector must be a bool scalar or vector value");
+ return false;
+ }
+
+ if (type->Is<type::Vector>()) {
+ auto size = type->As<type::Vector>()->size();
+ if (t2->is_scalar() || size != t2->As<type::Vector>()->size()) {
+ add_error(expr->params()[2]->source(),
"incorrect type for " + builtin +
- ". Input must be a runtime array");
+ ". Selector must be a vector with the same "
+ "number of elements as the result type");
+ return false;
+ }
+ } else {
+ if (!t2->is_scalar()) {
+ add_error(expr->params()[2]->source(),
+ "incorrect type for " + builtin +
+ ". Selector must be a bool scalar to match "
+ "scalar result type");
return false;
}
}
}
- // Result types don't match parameter types.
- if (data->intrinsic == ast::Intrinsic::kAll ||
- data->intrinsic == ast::Intrinsic::kAny) {
- if (!IsValidType(program_->TypeOf(expr->func()), expr->source(),
- builtin, IntrinsicDataType::kBoolScalar, 0,
- this)) {
+ if (data->intrinsic == semantic::Intrinsic::kArrayLength) {
+ if (!program_->TypeOf(expr)->UnwrapPtrIfNeeded()->Is<type::U32>()) {
+ add_error(
+ expr->source(),
+ "incorrect type for " + builtin +
+ ". Result type must be an unsigned int scalar value");
return false;
}
- }
- if (data->intrinsic == ast::Intrinsic::kDot) {
- if (!IsValidType(program_->TypeOf(expr->func()), expr->source(),
- builtin, IntrinsicDataType::kFloatScalar, 0,
- this)) {
- return false;
- }
- }
-
- if (data->intrinsic == ast::Intrinsic::kLength ||
- data->intrinsic == ast::Intrinsic::kDistance ||
- data->intrinsic == ast::Intrinsic::kDeterminant) {
- if (!IsValidType(program_->TypeOf(expr->func()), expr->source(),
- builtin, IntrinsicDataType::kFloatScalar, 0,
- this)) {
- return false;
- }
- }
-
- // Must be a square matrix.
- if (data->intrinsic == ast::Intrinsic::kDeterminant) {
- const auto* matrix =
- program_->TypeOf(expr->params()[0])->As<type::Matrix>();
- if (matrix->rows() != matrix->columns()) {
+ auto* p0 = program_->TypeOf(expr->params()[0])->UnwrapPtrIfNeeded();
+ if (!p0->Is<type::Array>() ||
+ !p0->As<type::Array>()->IsRuntimeArray()) {
add_error(expr->params()[0]->source(),
"incorrect type for " + builtin +
- ". Requires a square matrix");
+ ". Input must be a runtime array");
return false;
}
}
}
- // Last parameter must be a pointer.
- if (data->intrinsic == ast::Intrinsic::kFrexp ||
- data->intrinsic == ast::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 +
- ". Requires pointer value");
+ // Result types don't match parameter types.
+ if (data->intrinsic == semantic::Intrinsic::kAll ||
+ data->intrinsic == semantic::Intrinsic::kAny) {
+ if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
+ IntrinsicDataType::kBoolScalar, 0, this)) {
+ return false;
+ }
+ }
+
+ if (data->intrinsic == semantic::Intrinsic::kDot) {
+ if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
+ IntrinsicDataType::kFloatScalar, 0, this)) {
+ return false;
+ }
+ }
+
+ if (data->intrinsic == semantic::Intrinsic::kLength ||
+ data->intrinsic == semantic::Intrinsic::kDistance ||
+ data->intrinsic == semantic::Intrinsic::kDeterminant) {
+ if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
+ IntrinsicDataType::kFloatScalar, 0, this)) {
+ return false;
+ }
+ }
+
+ // Must be a square matrix.
+ 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(),
+ "incorrect type for " + builtin + ". Requires a square matrix");
return false;
}
}
}
- } else {
- if (!function_stack_.has(symbol)) {
- add_error(expr->source(), "v-0005",
- "function must be declared before use: '" +
- program_->Symbols().NameFor(symbol) + "'");
- return false;
- }
- if (symbol == current_function_->symbol()) {
- add_error(expr->source(), "v-0004",
- "recursion is not allowed: '" +
- program_->Symbols().NameFor(symbol) + "'");
- return false;
+
+ // Last parameter must be a pointer.
+ if (data->intrinsic == semantic::Intrinsic::kFrexp ||
+ 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 +
+ ". Requires pointer value");
+ return false;
+ }
}
}
+ 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: '" +
+ program_->Symbols().NameFor(symbol) + "'");
+ return false;
+ }
+ if (symbol == current_function_->symbol()) {
+ add_error(expr->source(), "v-0004",
+ "recursion is not allowed: '" +
+ program_->Symbols().NameFor(symbol) + "'");
+ return false;
+ }
+
} else {
add_error(expr->source(), "Invalid function call expression");
return false;
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 6e4eff6..fa6548d 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/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 @@
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())) {
- return EmitTextureCall(pre, out, expr);
- }
- auto name = generate_builtin_name(ident);
+ auto name = generate_builtin_name(sem);
if (name.empty()) {
return false;
}
@@ -634,30 +636,29 @@
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*>(
- ident->intrinsic_signature());
- auto& pidx = signature->params.idx;
- auto const kNotUsed = ast::intrinsic::TextureSignature::Parameters::kNotUsed;
+ auto& pidx = sem->Params().idx;
+ auto const kNotUsed = semantic::TextureIntrinsicCall::Parameters::kNotUsed;
auto* texture = params[pidx.texture];
auto* texture_type = TypeOf(texture)->UnwrapAll()->As<type::Texture>();
- switch (ident->intrinsic()) {
- case ast::Intrinsic::kTextureDimensions:
- case ast::Intrinsic::kTextureNumLayers:
- case ast::Intrinsic::kTextureNumLevels:
- case ast::Intrinsic::kTextureNumSamples: {
+ switch (sem->intrinsic()) {
+ case semantic::Intrinsic::kTextureDimensions:
+ case semantic::Intrinsic::kTextureNumLayers:
+ case semantic::Intrinsic::kTextureNumLevels:
+ 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()) {
- case ast::Intrinsic::kTextureDimensions:
+ switch (sem->intrinsic()) {
+ case semantic::Intrinsic::kTextureDimensions:
switch (texture_type->dim()) {
case type::TextureDimension::kNone:
error_ = "texture dimension is kNone";
@@ -693,7 +694,7 @@
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 @@
break;
}
break;
- case ast::Intrinsic::kTextureNumLevels:
+ case semantic::Intrinsic::kTextureNumLevels:
add_mip_level_in = true;
switch (texture_type->dim()) {
default:
@@ -728,7 +729,7 @@
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 pack_mip_in_coords = false;
- switch (ident->intrinsic()) {
- case ast::Intrinsic::kTextureSample:
+ switch (sem->intrinsic()) {
+ 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 @@
}
}
- 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 @@
} // namespace hlsl
std::string GeneratorImpl::generate_builtin_name(
- ast::IdentifierExpression* ident) {
+ const semantic::IntrinsicCall* call) {
std::string out;
- switch (ident->intrinsic()) {
- case ast::Intrinsic::kAcos:
- case ast::Intrinsic::kAny:
- case ast::Intrinsic::kAll:
- case ast::Intrinsic::kAsin:
- case ast::Intrinsic::kAtan:
- case ast::Intrinsic::kAtan2:
- case ast::Intrinsic::kCeil:
- case ast::Intrinsic::kCos:
- case ast::Intrinsic::kCosh:
- case ast::Intrinsic::kCross:
- case ast::Intrinsic::kDeterminant:
- case ast::Intrinsic::kDistance:
- case ast::Intrinsic::kDot:
- case ast::Intrinsic::kExp:
- case ast::Intrinsic::kExp2:
- case ast::Intrinsic::kFloor:
- case ast::Intrinsic::kFma:
- case ast::Intrinsic::kLdexp:
- case ast::Intrinsic::kLength:
- case ast::Intrinsic::kLog:
- case ast::Intrinsic::kLog2:
- case ast::Intrinsic::kNormalize:
- case ast::Intrinsic::kPow:
- case ast::Intrinsic::kReflect:
- case ast::Intrinsic::kRound:
- case ast::Intrinsic::kSin:
- case ast::Intrinsic::kSinh:
- case ast::Intrinsic::kSqrt:
- case ast::Intrinsic::kStep:
- case ast::Intrinsic::kTan:
- case ast::Intrinsic::kTanh:
- case ast::Intrinsic::kTrunc:
- case ast::Intrinsic::kMix:
- case ast::Intrinsic::kSign:
- case ast::Intrinsic::kAbs:
- case ast::Intrinsic::kMax:
- case ast::Intrinsic::kMin:
- case ast::Intrinsic::kClamp:
- out = builder_.Symbols().NameFor(ident->symbol());
+ switch (call->intrinsic()) {
+ case semantic::Intrinsic::kAcos:
+ case semantic::Intrinsic::kAny:
+ case semantic::Intrinsic::kAll:
+ case semantic::Intrinsic::kAsin:
+ case semantic::Intrinsic::kAtan:
+ case semantic::Intrinsic::kAtan2:
+ case semantic::Intrinsic::kCeil:
+ case semantic::Intrinsic::kCos:
+ case semantic::Intrinsic::kCosh:
+ case semantic::Intrinsic::kCross:
+ case semantic::Intrinsic::kDeterminant:
+ case semantic::Intrinsic::kDistance:
+ case semantic::Intrinsic::kDot:
+ case semantic::Intrinsic::kExp:
+ case semantic::Intrinsic::kExp2:
+ case semantic::Intrinsic::kFloor:
+ case semantic::Intrinsic::kFma:
+ case semantic::Intrinsic::kLdexp:
+ case semantic::Intrinsic::kLength:
+ case semantic::Intrinsic::kLog:
+ case semantic::Intrinsic::kLog2:
+ case semantic::Intrinsic::kNormalize:
+ case semantic::Intrinsic::kPow:
+ case semantic::Intrinsic::kReflect:
+ case semantic::Intrinsic::kRound:
+ case semantic::Intrinsic::kSin:
+ case semantic::Intrinsic::kSinh:
+ case semantic::Intrinsic::kSqrt:
+ case semantic::Intrinsic::kStep:
+ case semantic::Intrinsic::kTan:
+ case semantic::Intrinsic::kTanh:
+ case semantic::Intrinsic::kTrunc:
+ case semantic::Intrinsic::kMix:
+ case semantic::Intrinsic::kSign:
+ case semantic::Intrinsic::kAbs:
+ case semantic::Intrinsic::kMax:
+ case semantic::Intrinsic::kMin:
+ case semantic::Intrinsic::kClamp:
+ 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 ast::Intrinsic::kFwidthCoarse:
- case ast::Intrinsic::kFwidthFine:
+ case semantic::Intrinsic::kFwidth:
+ case semantic::Intrinsic::kFwidthCoarse:
+ 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 "";
}
diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h
index 0c175fc..acd0ae1 100644
--- a/src/writer/hlsl/generator_impl.h
+++ b/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 @@
/// @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 @@
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
diff --git a/src/writer/hlsl/generator_impl_binary_test.cc b/src/writer/hlsl/generator_impl_binary_test.cc
index 379259e..bbe1eca 100644
--- a/src/writer/hlsl/generator_impl_binary_test.cc
+++ b/src/writer/hlsl/generator_impl_binary_test.cc
@@ -483,6 +483,10 @@
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 @@
Expr("d"))));
auto* expr = create<ast::CallStatement>(Call("foo", params));
+ WrapInFunction(expr);
GeneratorImpl& gen = Build();
diff --git a/src/writer/hlsl/generator_impl_call_test.cc b/src/writer/hlsl/generator_impl_call_test.cc
index bf7661f..8094390 100644
--- a/src/writer/hlsl/generator_impl_call_test.cc
+++ b/src/writer/hlsl/generator_impl_call_test.cc
@@ -30,11 +30,12 @@
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_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, 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();
diff --git a/src/writer/hlsl/generator_impl_intrinsic_test.cc b/src/writer/hlsl/generator_impl_intrinsic_test.cc
index a7aadbf..568ecb0 100644
--- a/src/writer/hlsl/generator_impl_intrinsic_test.cc
+++ b/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 @@
};
struct IntrinsicData {
- ast::Intrinsic intrinsic;
+ semantic::Intrinsic intrinsic;
ParamType type;
const char* hlsl_name;
};
@@ -57,92 +58,92 @@
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 ast::Intrinsic::kAsin:
- case ast::Intrinsic::kAtan:
- case ast::Intrinsic::kCeil:
- case ast::Intrinsic::kCos:
- case ast::Intrinsic::kCosh:
- case ast::Intrinsic::kDpdx:
- case ast::Intrinsic::kDpdxCoarse:
- case ast::Intrinsic::kDpdxFine:
- case ast::Intrinsic::kDpdy:
- case ast::Intrinsic::kDpdyCoarse:
- case ast::Intrinsic::kDpdyFine:
- case ast::Intrinsic::kExp:
- case ast::Intrinsic::kExp2:
- case ast::Intrinsic::kFloor:
- case ast::Intrinsic::kFract:
- case ast::Intrinsic::kFwidth:
- case ast::Intrinsic::kFwidthCoarse:
- case ast::Intrinsic::kFwidthFine:
- case ast::Intrinsic::kInverseSqrt:
- case ast::Intrinsic::kIsFinite:
- case ast::Intrinsic::kIsInf:
- case ast::Intrinsic::kIsNan:
- case ast::Intrinsic::kIsNormal:
- case ast::Intrinsic::kLdexp:
- case ast::Intrinsic::kLength:
- case ast::Intrinsic::kLog:
- case ast::Intrinsic::kLog2:
- case ast::Intrinsic::kNormalize:
- case ast::Intrinsic::kReflect:
- case ast::Intrinsic::kRound:
- case ast::Intrinsic::kSin:
- case ast::Intrinsic::kSinh:
- case ast::Intrinsic::kSqrt:
- case ast::Intrinsic::kTan:
- case ast::Intrinsic::kTanh:
- case ast::Intrinsic::kTrunc:
- case ast::Intrinsic::kSign:
+ case semantic::Intrinsic::kAcos:
+ case semantic::Intrinsic::kAsin:
+ case semantic::Intrinsic::kAtan:
+ case semantic::Intrinsic::kCeil:
+ case semantic::Intrinsic::kCos:
+ case semantic::Intrinsic::kCosh:
+ case semantic::Intrinsic::kDpdx:
+ case semantic::Intrinsic::kDpdxCoarse:
+ case semantic::Intrinsic::kDpdxFine:
+ case semantic::Intrinsic::kDpdy:
+ case semantic::Intrinsic::kDpdyCoarse:
+ case semantic::Intrinsic::kDpdyFine:
+ case semantic::Intrinsic::kExp:
+ case semantic::Intrinsic::kExp2:
+ case semantic::Intrinsic::kFloor:
+ case semantic::Intrinsic::kFract:
+ case semantic::Intrinsic::kFwidth:
+ case semantic::Intrinsic::kFwidthCoarse:
+ case semantic::Intrinsic::kFwidthFine:
+ case semantic::Intrinsic::kInverseSqrt:
+ case semantic::Intrinsic::kIsFinite:
+ case semantic::Intrinsic::kIsInf:
+ case semantic::Intrinsic::kIsNan:
+ case semantic::Intrinsic::kIsNormal:
+ case semantic::Intrinsic::kLdexp:
+ case semantic::Intrinsic::kLength:
+ case semantic::Intrinsic::kLog:
+ case semantic::Intrinsic::kLog2:
+ case semantic::Intrinsic::kNormalize:
+ case semantic::Intrinsic::kReflect:
+ case semantic::Intrinsic::kRound:
+ case semantic::Intrinsic::kSin:
+ case semantic::Intrinsic::kSinh:
+ case semantic::Intrinsic::kSqrt:
+ case semantic::Intrinsic::kTan:
+ case semantic::Intrinsic::kTanh:
+ case semantic::Intrinsic::kTrunc:
+ case semantic::Intrinsic::kSign:
return builder->Call(str.str(), "f1");
- case ast::Intrinsic::kAtan2:
- case ast::Intrinsic::kCross:
- case ast::Intrinsic::kDot:
- case ast::Intrinsic::kDistance:
- case ast::Intrinsic::kPow:
- case ast::Intrinsic::kStep:
+ case semantic::Intrinsic::kAtan2:
+ case semantic::Intrinsic::kCross:
+ case semantic::Intrinsic::kDot:
+ case semantic::Intrinsic::kDistance:
+ case semantic::Intrinsic::kPow:
+ case semantic::Intrinsic::kStep:
return builder->Call(str.str(), "f1", "f2");
- case ast::Intrinsic::kFma:
- case ast::Intrinsic::kMix:
- case ast::Intrinsic::kFaceForward:
- case ast::Intrinsic::kSmoothStep:
+ case semantic::Intrinsic::kFma:
+ case semantic::Intrinsic::kMix:
+ case semantic::Intrinsic::kFaceForward:
+ case semantic::Intrinsic::kSmoothStep:
return builder->Call(str.str(), "f1", "f2", "f3");
- case ast::Intrinsic::kAll:
- case ast::Intrinsic::kAny:
+ case semantic::Intrinsic::kAll:
+ 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 ast::Intrinsic::kReverseBits:
+ case semantic::Intrinsic::kCountOneBits:
+ case semantic::Intrinsic::kReverseBits:
return builder->Call(str.str(), "u1");
- case ast::Intrinsic::kMax:
- case ast::Intrinsic::kMin:
+ case semantic::Intrinsic::kMax:
+ 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 @@
GeneratorImpl& gen = Build();
- EXPECT_EQ(
- gen.generate_builtin_name(call->func()->As<ast::IdentifierExpression>()),
- param.hlsl_name);
+ auto* sem = program->Sem().Get(call);
+ ASSERT_NE(sem, nullptr);
+ 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{ast::Intrinsic::kAbs, ParamType::kU32, "abs"},
- IntrinsicData{ast::Intrinsic::kAcos, ParamType::kF32, "acos"},
- IntrinsicData{ast::Intrinsic::kAll, ParamType::kBool, "all"},
- IntrinsicData{ast::Intrinsic::kAny, ParamType::kBool, "any"},
- IntrinsicData{ast::Intrinsic::kAsin, ParamType::kF32, "asin"},
- IntrinsicData{ast::Intrinsic::kAtan, ParamType::kF32, "atan"},
- IntrinsicData{ast::Intrinsic::kAtan2, ParamType::kF32, "atan2"},
- IntrinsicData{ast::Intrinsic::kCeil, ParamType::kF32, "ceil"},
- IntrinsicData{ast::Intrinsic::kClamp, ParamType::kF32, "clamp"},
- IntrinsicData{ast::Intrinsic::kClamp, ParamType::kU32, "clamp"},
- IntrinsicData{ast::Intrinsic::kCos, ParamType::kF32, "cos"},
- IntrinsicData{ast::Intrinsic::kCosh, ParamType::kF32, "cosh"},
- IntrinsicData{ast::Intrinsic::kCountOneBits, ParamType::kU32,
+ IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kF32, "abs"},
+ IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kU32, "abs"},
+ IntrinsicData{semantic::Intrinsic::kAcos, ParamType::kF32, "acos"},
+ IntrinsicData{semantic::Intrinsic::kAll, ParamType::kBool, "all"},
+ IntrinsicData{semantic::Intrinsic::kAny, ParamType::kBool, "any"},
+ IntrinsicData{semantic::Intrinsic::kAsin, ParamType::kF32, "asin"},
+ IntrinsicData{semantic::Intrinsic::kAtan, ParamType::kF32, "atan"},
+ IntrinsicData{semantic::Intrinsic::kAtan2, ParamType::kF32, "atan2"},
+ IntrinsicData{semantic::Intrinsic::kCeil, ParamType::kF32, "ceil"},
+ IntrinsicData{semantic::Intrinsic::kClamp, ParamType::kF32, "clamp"},
+ IntrinsicData{semantic::Intrinsic::kClamp, ParamType::kU32, "clamp"},
+ IntrinsicData{semantic::Intrinsic::kCos, ParamType::kF32, "cos"},
+ IntrinsicData{semantic::Intrinsic::kCosh, ParamType::kF32, "cosh"},
+ IntrinsicData{semantic::Intrinsic::kCountOneBits, ParamType::kU32,
"countbits"},
- IntrinsicData{ast::Intrinsic::kCross, ParamType::kF32, "cross"},
- IntrinsicData{ast::Intrinsic::kDeterminant, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kCross, ParamType::kF32, "cross"},
+ IntrinsicData{semantic::Intrinsic::kDeterminant, ParamType::kF32,
"determinant"},
- IntrinsicData{ast::Intrinsic::kDistance, ParamType::kF32, "distance"},
- IntrinsicData{ast::Intrinsic::kDot, ParamType::kF32, "dot"},
- IntrinsicData{ast::Intrinsic::kDpdx, ParamType::kF32, "ddx"},
- IntrinsicData{ast::Intrinsic::kDpdxCoarse, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kDistance, ParamType::kF32,
+ "distance"},
+ IntrinsicData{semantic::Intrinsic::kDot, ParamType::kF32, "dot"},
+ IntrinsicData{semantic::Intrinsic::kDpdx, ParamType::kF32, "ddx"},
+ IntrinsicData{semantic::Intrinsic::kDpdxCoarse, ParamType::kF32,
"ddx_coarse"},
- IntrinsicData{ast::Intrinsic::kDpdxFine, ParamType::kF32, "ddx_fine"},
- IntrinsicData{ast::Intrinsic::kDpdy, ParamType::kF32, "ddy"},
- IntrinsicData{ast::Intrinsic::kDpdyCoarse, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kDpdxFine, ParamType::kF32,
+ "ddx_fine"},
+ IntrinsicData{semantic::Intrinsic::kDpdy, ParamType::kF32, "ddy"},
+ IntrinsicData{semantic::Intrinsic::kDpdyCoarse, ParamType::kF32,
"ddy_coarse"},
- IntrinsicData{ast::Intrinsic::kDpdyFine, ParamType::kF32, "ddy_fine"},
- IntrinsicData{ast::Intrinsic::kExp, ParamType::kF32, "exp"},
- IntrinsicData{ast::Intrinsic::kExp2, ParamType::kF32, "exp2"},
- IntrinsicData{ast::Intrinsic::kFaceForward, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kDpdyFine, ParamType::kF32,
+ "ddy_fine"},
+ IntrinsicData{semantic::Intrinsic::kExp, ParamType::kF32, "exp"},
+ IntrinsicData{semantic::Intrinsic::kExp2, ParamType::kF32, "exp2"},
+ IntrinsicData{semantic::Intrinsic::kFaceForward, ParamType::kF32,
"faceforward"},
- IntrinsicData{ast::Intrinsic::kFloor, ParamType::kF32, "floor"},
- IntrinsicData{ast::Intrinsic::kFma, ParamType::kF32, "fma"},
- IntrinsicData{ast::Intrinsic::kFract, ParamType::kF32, "frac"},
- IntrinsicData{ast::Intrinsic::kFwidth, ParamType::kF32, "fwidth"},
- IntrinsicData{ast::Intrinsic::kFwidthCoarse, ParamType::kF32, "fwidth"},
- IntrinsicData{ast::Intrinsic::kFwidthFine, ParamType::kF32, "fwidth"},
- IntrinsicData{ast::Intrinsic::kInverseSqrt, ParamType::kF32, "rsqrt"},
- IntrinsicData{ast::Intrinsic::kIsFinite, ParamType::kF32, "isfinite"},
- IntrinsicData{ast::Intrinsic::kIsInf, ParamType::kF32, "isinf"},
- IntrinsicData{ast::Intrinsic::kIsNan, ParamType::kF32, "isnan"},
- IntrinsicData{ast::Intrinsic::kLdexp, ParamType::kF32, "ldexp"},
- IntrinsicData{ast::Intrinsic::kLength, ParamType::kF32, "length"},
- IntrinsicData{ast::Intrinsic::kLog, ParamType::kF32, "log"},
- IntrinsicData{ast::Intrinsic::kLog2, ParamType::kF32, "log2"},
- IntrinsicData{ast::Intrinsic::kMax, ParamType::kF32, "max"},
- IntrinsicData{ast::Intrinsic::kMax, ParamType::kU32, "max"},
- IntrinsicData{ast::Intrinsic::kMin, ParamType::kF32, "min"},
- IntrinsicData{ast::Intrinsic::kMin, ParamType::kU32, "min"},
- IntrinsicData{ast::Intrinsic::kNormalize, ParamType::kF32, "normalize"},
- IntrinsicData{ast::Intrinsic::kPow, ParamType::kF32, "pow"},
- IntrinsicData{ast::Intrinsic::kReflect, ParamType::kF32, "reflect"},
- IntrinsicData{ast::Intrinsic::kReverseBits, ParamType::kU32,
+ IntrinsicData{semantic::Intrinsic::kFloor, ParamType::kF32, "floor"},
+ IntrinsicData{semantic::Intrinsic::kFma, ParamType::kF32, "fma"},
+ IntrinsicData{semantic::Intrinsic::kFract, ParamType::kF32, "frac"},
+ IntrinsicData{semantic::Intrinsic::kFwidth, ParamType::kF32, "fwidth"},
+ IntrinsicData{semantic::Intrinsic::kFwidthCoarse, ParamType::kF32,
+ "fwidth"},
+ IntrinsicData{semantic::Intrinsic::kFwidthFine, ParamType::kF32,
+ "fwidth"},
+ IntrinsicData{semantic::Intrinsic::kInverseSqrt, ParamType::kF32,
+ "rsqrt"},
+ IntrinsicData{semantic::Intrinsic::kIsFinite, ParamType::kF32,
+ "isfinite"},
+ IntrinsicData{semantic::Intrinsic::kIsInf, ParamType::kF32, "isinf"},
+ IntrinsicData{semantic::Intrinsic::kIsNan, ParamType::kF32, "isnan"},
+ IntrinsicData{semantic::Intrinsic::kLdexp, ParamType::kF32, "ldexp"},
+ IntrinsicData{semantic::Intrinsic::kLength, ParamType::kF32, "length"},
+ IntrinsicData{semantic::Intrinsic::kLog, ParamType::kF32, "log"},
+ IntrinsicData{semantic::Intrinsic::kLog2, ParamType::kF32, "log2"},
+ IntrinsicData{semantic::Intrinsic::kMax, ParamType::kF32, "max"},
+ IntrinsicData{semantic::Intrinsic::kMax, ParamType::kU32, "max"},
+ IntrinsicData{semantic::Intrinsic::kMin, ParamType::kF32, "min"},
+ 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{ast::Intrinsic::kSign, ParamType::kF32, "sign"},
- IntrinsicData{ast::Intrinsic::kSin, ParamType::kF32, "sin"},
- IntrinsicData{ast::Intrinsic::kSinh, ParamType::kF32, "sinh"},
- IntrinsicData{ast::Intrinsic::kSmoothStep, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kRound, ParamType::kU32, "round"},
+ IntrinsicData{semantic::Intrinsic::kSign, ParamType::kF32, "sign"},
+ IntrinsicData{semantic::Intrinsic::kSin, ParamType::kF32, "sin"},
+ IntrinsicData{semantic::Intrinsic::kSinh, ParamType::kF32, "sinh"},
+ IntrinsicData{semantic::Intrinsic::kSmoothStep, ParamType::kF32,
"smoothstep"},
- IntrinsicData{ast::Intrinsic::kSqrt, ParamType::kF32, "sqrt"},
- IntrinsicData{ast::Intrinsic::kStep, ParamType::kF32, "step"},
- IntrinsicData{ast::Intrinsic::kTan, ParamType::kF32, "tan"},
- IntrinsicData{ast::Intrinsic::kTanh, ParamType::kF32, "tanh"},
- IntrinsicData{ast::Intrinsic::kTrunc, ParamType::kF32, "trunc"}));
+ IntrinsicData{semantic::Intrinsic::kSqrt, ParamType::kF32, "sqrt"},
+ IntrinsicData{semantic::Intrinsic::kStep, ParamType::kF32, "step"},
+ IntrinsicData{semantic::Intrinsic::kTan, ParamType::kF32, "tan"},
+ IntrinsicData{semantic::Intrinsic::kTanh, ParamType::kF32, "tanh"},
+ IntrinsicData{semantic::Intrinsic::kTrunc, ParamType::kF32, "trunc"}));
TEST_F(HlslGeneratorImplTest_Intrinsic, DISABLED_Intrinsic_IsNormal) {
FAIL();
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index be67f4d..77678d8 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/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 @@
return 0;
}
- if (ident->IsIntrinsic()) {
- if (ast::intrinsic::IsTextureIntrinsic(ident->intrinsic())) {
- return EmitTextureCall(expr);
- }
- auto name = generate_builtin_name(ident);
+ auto* call_sem = program_->Sem().Get(expr);
+ if (auto* sem = call_sem->As<semantic::TextureIntrinsicCall>()) {
+ return EmitTextureCall(expr, sem);
+ }
+ if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) {
+ auto name = generate_builtin_name(sem);
if (name.empty()) {
return false;
}
@@ -555,21 +557,20 @@
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*>(
- ident->intrinsic_signature());
- auto& pidx = signature->params.idx;
- auto const kNotUsed = ast::intrinsic::TextureSignature::Parameters::kNotUsed;
+ auto& pidx = sem->Params().idx;
+ auto const kNotUsed = semantic::TextureIntrinsicCall::Parameters::kNotUsed;
assert(pidx.texture != kNotUsed);
auto* texture_type =
TypeOf(params[pidx.texture])->UnwrapAll()->As<type::Texture>();
- switch (ident->intrinsic()) {
- case ast::Intrinsic::kTextureDimensions: {
+ switch (sem->intrinsic()) {
+ case semantic::Intrinsic::kTextureDimensions: {
std::vector<const char*> dims;
switch (texture_type->dim()) {
case type::TextureDimension::kNone:
@@ -623,7 +624,7 @@
}
return true;
}
- case ast::Intrinsic::kTextureNumLayers: {
+ case semantic::Intrinsic::kTextureNumLayers: {
out_ << "int(";
if (!EmitExpression(params[pidx.texture])) {
return false;
@@ -631,7 +632,7 @@
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 @@
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 lod_param_is_named = true;
- switch (ident->intrinsic()) {
- case ast::Intrinsic::kTextureSample:
- case ast::Intrinsic::kTextureSampleBias:
- case ast::Intrinsic::kTextureSampleLevel:
- case ast::Intrinsic::kTextureSampleGrad:
+ switch (sem->intrinsic()) {
+ case semantic::Intrinsic::kTextureSample:
+ case semantic::Intrinsic::kTextureSampleBias:
+ case semantic::Intrinsic::kTextureSampleLevel:
+ 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 @@
}
std::string GeneratorImpl::generate_builtin_name(
- ast::IdentifierExpression* ident) {
- auto* type = TypeOf(ident);
+ const semantic::IntrinsicCall* call) {
std::string out = "metal::";
- switch (ident->intrinsic()) {
- case ast::Intrinsic::kAcos:
- case ast::Intrinsic::kAll:
- case ast::Intrinsic::kAny:
- case ast::Intrinsic::kAsin:
- case ast::Intrinsic::kAtan:
- case ast::Intrinsic::kAtan2:
- case ast::Intrinsic::kCeil:
- case ast::Intrinsic::kCos:
- case ast::Intrinsic::kCosh:
- case ast::Intrinsic::kCross:
- case ast::Intrinsic::kDeterminant:
- case ast::Intrinsic::kDistance:
- case ast::Intrinsic::kDot:
- case ast::Intrinsic::kExp:
- case ast::Intrinsic::kExp2:
- case ast::Intrinsic::kFloor:
- case ast::Intrinsic::kFma:
- case ast::Intrinsic::kFract:
- case ast::Intrinsic::kLength:
- case ast::Intrinsic::kLdexp:
- case ast::Intrinsic::kLog:
- case ast::Intrinsic::kLog2:
- case ast::Intrinsic::kMix:
- case ast::Intrinsic::kNormalize:
- case ast::Intrinsic::kPow:
- case ast::Intrinsic::kReflect:
- case ast::Intrinsic::kRound:
- case ast::Intrinsic::kSelect:
- case ast::Intrinsic::kSin:
- case ast::Intrinsic::kSinh:
- case ast::Intrinsic::kSqrt:
- case ast::Intrinsic::kStep:
- case ast::Intrinsic::kTan:
- case ast::Intrinsic::kTanh:
- case ast::Intrinsic::kTrunc:
- case ast::Intrinsic::kSign:
- case ast::Intrinsic::kClamp:
- out += program_->Symbols().NameFor(ident->symbol());
+ switch (call->intrinsic()) {
+ case semantic::Intrinsic::kAcos:
+ case semantic::Intrinsic::kAll:
+ case semantic::Intrinsic::kAny:
+ case semantic::Intrinsic::kAsin:
+ case semantic::Intrinsic::kAtan:
+ case semantic::Intrinsic::kAtan2:
+ case semantic::Intrinsic::kCeil:
+ case semantic::Intrinsic::kCos:
+ case semantic::Intrinsic::kCosh:
+ case semantic::Intrinsic::kCross:
+ case semantic::Intrinsic::kDeterminant:
+ case semantic::Intrinsic::kDistance:
+ case semantic::Intrinsic::kDot:
+ case semantic::Intrinsic::kExp:
+ case semantic::Intrinsic::kExp2:
+ case semantic::Intrinsic::kFloor:
+ case semantic::Intrinsic::kFma:
+ case semantic::Intrinsic::kFract:
+ case semantic::Intrinsic::kLength:
+ case semantic::Intrinsic::kLdexp:
+ case semantic::Intrinsic::kLog:
+ case semantic::Intrinsic::kLog2:
+ case semantic::Intrinsic::kMix:
+ case semantic::Intrinsic::kNormalize:
+ case semantic::Intrinsic::kPow:
+ case semantic::Intrinsic::kReflect:
+ case semantic::Intrinsic::kRound:
+ case semantic::Intrinsic::kSelect:
+ case semantic::Intrinsic::kSin:
+ case semantic::Intrinsic::kSinh:
+ case semantic::Intrinsic::kSqrt:
+ case semantic::Intrinsic::kStep:
+ case semantic::Intrinsic::kTan:
+ case semantic::Intrinsic::kTanh:
+ case semantic::Intrinsic::kTrunc:
+ case semantic::Intrinsic::kSign:
+ case semantic::Intrinsic::kClamp:
+ out += semantic::intrinsic::str(call->intrinsic());
break;
- case ast::Intrinsic::kAbs:
- if (type->is_float_scalar_or_vector()) {
+ case semantic::Intrinsic::kAbs:
+ 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 ast::Intrinsic::kDpdxCoarse:
- case ast::Intrinsic::kDpdxFine:
+ case semantic::Intrinsic::kDpdx:
+ case semantic::Intrinsic::kDpdxCoarse:
+ case semantic::Intrinsic::kDpdxFine:
out += "dfdx";
break;
- case ast::Intrinsic::kDpdy:
- case ast::Intrinsic::kDpdyCoarse:
- case ast::Intrinsic::kDpdyFine:
+ case semantic::Intrinsic::kDpdy:
+ case semantic::Intrinsic::kDpdyCoarse:
+ case semantic::Intrinsic::kDpdyFine:
out += "dfdy";
break;
- case ast::Intrinsic::kFwidth:
- case ast::Intrinsic::kFwidthCoarse:
- case ast::Intrinsic::kFwidthFine:
+ case semantic::Intrinsic::kFwidth:
+ case semantic::Intrinsic::kFwidthCoarse:
+ 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:
- if (type->is_float_scalar_or_vector()) {
+ case semantic::Intrinsic::kMax:
+ if (call->Type()->is_float_scalar_or_vector()) {
out += "fmax";
} else {
out += "max";
}
break;
- case ast::Intrinsic::kMin:
- if (type->is_float_scalar_or_vector()) {
+ case semantic::Intrinsic::kMin:
+ 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;
diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h
index 0aafa62..617061f 100644
--- a/src/writer/msl/generator_impl.h
+++ b/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 @@
/// 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 @@
/// @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
diff --git a/src/writer/msl/generator_impl_call_test.cc b/src/writer/msl/generator_impl_call_test.cc
index a031f3f..172f402 100644
--- a/src/writer/msl/generator_impl_call_test.cc
+++ b/src/writer/msl/generator_impl_call_test.cc
@@ -32,10 +32,12 @@
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_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, 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");
}
diff --git a/src/writer/msl/generator_impl_import_test.cc b/src/writer/msl/generator_impl_import_test.cc
index 7700223..76440d0 100644
--- a/src/writer/msl/generator_impl_import_test.cc
+++ b/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,9 +58,13 @@
GeneratorImpl& gen = Build();
- ASSERT_EQ(
- gen.generate_builtin_name(call->func()->As<ast::IdentifierExpression>()),
- std::string("metal::") + param.msl_name);
+ auto* sem = program->Sem().Get(call);
+ 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,
MslImportData_SingleParamTest,
diff --git a/src/writer/msl/generator_impl_intrinsic_test.cc b/src/writer/msl/generator_impl_intrinsic_test.cc
index 133c196..bd56c6f 100644
--- a/src/writer/msl/generator_impl_intrinsic_test.cc
+++ b/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 @@
};
struct IntrinsicData {
- ast::Intrinsic intrinsic;
+ semantic::Intrinsic intrinsic;
ParamType type;
const char* msl_name;
};
@@ -59,92 +60,92 @@
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 ast::Intrinsic::kAsin:
- case ast::Intrinsic::kAtan:
- case ast::Intrinsic::kCeil:
- case ast::Intrinsic::kCos:
- case ast::Intrinsic::kCosh:
- case ast::Intrinsic::kDpdx:
- case ast::Intrinsic::kDpdxCoarse:
- case ast::Intrinsic::kDpdxFine:
- case ast::Intrinsic::kDpdy:
- case ast::Intrinsic::kDpdyCoarse:
- case ast::Intrinsic::kDpdyFine:
- case ast::Intrinsic::kExp:
- case ast::Intrinsic::kExp2:
- case ast::Intrinsic::kFloor:
- case ast::Intrinsic::kFract:
- case ast::Intrinsic::kFwidth:
- case ast::Intrinsic::kFwidthCoarse:
- case ast::Intrinsic::kFwidthFine:
- case ast::Intrinsic::kInverseSqrt:
- case ast::Intrinsic::kIsFinite:
- case ast::Intrinsic::kIsInf:
- case ast::Intrinsic::kIsNan:
- case ast::Intrinsic::kIsNormal:
- case ast::Intrinsic::kLdexp:
- case ast::Intrinsic::kLength:
- case ast::Intrinsic::kLog:
- case ast::Intrinsic::kLog2:
- case ast::Intrinsic::kNormalize:
- case ast::Intrinsic::kReflect:
- case ast::Intrinsic::kRound:
- case ast::Intrinsic::kSin:
- case ast::Intrinsic::kSinh:
- case ast::Intrinsic::kSqrt:
- case ast::Intrinsic::kTan:
- case ast::Intrinsic::kTanh:
- case ast::Intrinsic::kTrunc:
- case ast::Intrinsic::kSign:
+ case semantic::Intrinsic::kAcos:
+ case semantic::Intrinsic::kAsin:
+ case semantic::Intrinsic::kAtan:
+ case semantic::Intrinsic::kCeil:
+ case semantic::Intrinsic::kCos:
+ case semantic::Intrinsic::kCosh:
+ case semantic::Intrinsic::kDpdx:
+ case semantic::Intrinsic::kDpdxCoarse:
+ case semantic::Intrinsic::kDpdxFine:
+ case semantic::Intrinsic::kDpdy:
+ case semantic::Intrinsic::kDpdyCoarse:
+ case semantic::Intrinsic::kDpdyFine:
+ case semantic::Intrinsic::kExp:
+ case semantic::Intrinsic::kExp2:
+ case semantic::Intrinsic::kFloor:
+ case semantic::Intrinsic::kFract:
+ case semantic::Intrinsic::kFwidth:
+ case semantic::Intrinsic::kFwidthCoarse:
+ case semantic::Intrinsic::kFwidthFine:
+ case semantic::Intrinsic::kInverseSqrt:
+ case semantic::Intrinsic::kIsFinite:
+ case semantic::Intrinsic::kIsInf:
+ case semantic::Intrinsic::kIsNan:
+ case semantic::Intrinsic::kIsNormal:
+ case semantic::Intrinsic::kLdexp:
+ case semantic::Intrinsic::kLength:
+ case semantic::Intrinsic::kLog:
+ case semantic::Intrinsic::kLog2:
+ case semantic::Intrinsic::kNormalize:
+ case semantic::Intrinsic::kReflect:
+ case semantic::Intrinsic::kRound:
+ case semantic::Intrinsic::kSin:
+ case semantic::Intrinsic::kSinh:
+ case semantic::Intrinsic::kSqrt:
+ case semantic::Intrinsic::kTan:
+ case semantic::Intrinsic::kTanh:
+ case semantic::Intrinsic::kTrunc:
+ case semantic::Intrinsic::kSign:
return builder->Call(str.str(), "f1");
- case ast::Intrinsic::kAtan2:
- case ast::Intrinsic::kCross:
- case ast::Intrinsic::kDot:
- case ast::Intrinsic::kDistance:
- case ast::Intrinsic::kPow:
- case ast::Intrinsic::kStep:
+ case semantic::Intrinsic::kAtan2:
+ case semantic::Intrinsic::kCross:
+ case semantic::Intrinsic::kDot:
+ case semantic::Intrinsic::kDistance:
+ case semantic::Intrinsic::kPow:
+ case semantic::Intrinsic::kStep:
return builder->Call(str.str(), "f1", "f2");
- case ast::Intrinsic::kFma:
- case ast::Intrinsic::kMix:
- case ast::Intrinsic::kFaceForward:
- case ast::Intrinsic::kSmoothStep:
+ case semantic::Intrinsic::kFma:
+ case semantic::Intrinsic::kMix:
+ case semantic::Intrinsic::kFaceForward:
+ case semantic::Intrinsic::kSmoothStep:
return builder->Call(str.str(), "f1", "f2", "f3");
- case ast::Intrinsic::kAll:
- case ast::Intrinsic::kAny:
+ case semantic::Intrinsic::kAll:
+ 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 ast::Intrinsic::kReverseBits:
+ case semantic::Intrinsic::kCountOneBits:
+ case semantic::Intrinsic::kReverseBits:
return builder->Call(str.str(), "u1");
- case ast::Intrinsic::kMax:
- case ast::Intrinsic::kMin:
+ case semantic::Intrinsic::kMax:
+ 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 @@
GeneratorImpl& gen = Build();
- EXPECT_EQ(
- gen.generate_builtin_name(call->func()->As<ast::IdentifierExpression>()),
- param.msl_name);
+ auto* sem = program->Sem().Get(call);
+ ASSERT_NE(sem, nullptr);
+ 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{ast::Intrinsic::kAbs, ParamType::kU32, "metal::abs"},
- IntrinsicData{ast::Intrinsic::kAcos, ParamType::kF32, "metal::acos"},
- IntrinsicData{ast::Intrinsic::kAll, ParamType::kBool, "metal::all"},
- IntrinsicData{ast::Intrinsic::kAny, ParamType::kBool, "metal::any"},
- IntrinsicData{ast::Intrinsic::kAsin, ParamType::kF32, "metal::asin"},
- IntrinsicData{ast::Intrinsic::kAtan, ParamType::kF32, "metal::atan"},
- IntrinsicData{ast::Intrinsic::kAtan2, ParamType::kF32, "metal::atan2"},
- IntrinsicData{ast::Intrinsic::kCeil, ParamType::kF32, "metal::ceil"},
- IntrinsicData{ast::Intrinsic::kClamp, ParamType::kF32, "metal::clamp"},
- IntrinsicData{ast::Intrinsic::kClamp, ParamType::kU32, "metal::clamp"},
- IntrinsicData{ast::Intrinsic::kCos, ParamType::kF32, "metal::cos"},
- IntrinsicData{ast::Intrinsic::kCosh, ParamType::kF32, "metal::cosh"},
- IntrinsicData{ast::Intrinsic::kCountOneBits, ParamType::kU32,
+ IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kF32,
+ "metal::fabs"},
+ IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kU32, "metal::abs"},
+ IntrinsicData{semantic::Intrinsic::kAcos, ParamType::kF32,
+ "metal::acos"},
+ IntrinsicData{semantic::Intrinsic::kAll, ParamType::kBool,
+ "metal::all"},
+ IntrinsicData{semantic::Intrinsic::kAny, ParamType::kBool,
+ "metal::any"},
+ IntrinsicData{semantic::Intrinsic::kAsin, ParamType::kF32,
+ "metal::asin"},
+ IntrinsicData{semantic::Intrinsic::kAtan, ParamType::kF32,
+ "metal::atan"},
+ 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{ast::Intrinsic::kDeterminant, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kCross, 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{ast::Intrinsic::kDpdx, ParamType::kF32, "metal::dfdx"},
- IntrinsicData{ast::Intrinsic::kDpdxCoarse, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kDot, ParamType::kF32, "metal::dot"},
+ IntrinsicData{semantic::Intrinsic::kDpdx, 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{ast::Intrinsic::kDpdyCoarse, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kDpdxFine, 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{ast::Intrinsic::kExp2, ParamType::kF32, "metal::exp2"},
- IntrinsicData{ast::Intrinsic::kFaceForward, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kDpdyFine, ParamType::kF32,
+ "metal::dfdy"},
+ 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{ast::Intrinsic::kFma, ParamType::kF32, "metal::fma"},
- IntrinsicData{ast::Intrinsic::kFract, ParamType::kF32, "metal::fract"},
- IntrinsicData{ast::Intrinsic::kFwidth, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kFloor, ParamType::kF32,
+ "metal::floor"},
+ IntrinsicData{semantic::Intrinsic::kFma, ParamType::kF32, "metal::fma"},
+ 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{ast::Intrinsic::kIsNan, ParamType::kF32, "metal::isnan"},
- IntrinsicData{ast::Intrinsic::kIsNormal, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kIsInf, ParamType::kF32,
+ "metal::isinf"},
+ 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{ast::Intrinsic::kLength, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kLdexp, ParamType::kF32,
+ "metal::ldexp"},
+ IntrinsicData{semantic::Intrinsic::kLength, ParamType::kF32,
"metal::length"},
- IntrinsicData{ast::Intrinsic::kLog, ParamType::kF32, "metal::log"},
- IntrinsicData{ast::Intrinsic::kLog2, ParamType::kF32, "metal::log2"},
- IntrinsicData{ast::Intrinsic::kMax, ParamType::kF32, "metal::fmax"},
- IntrinsicData{ast::Intrinsic::kMax, ParamType::kU32, "metal::max"},
- IntrinsicData{ast::Intrinsic::kMin, ParamType::kF32, "metal::fmin"},
- IntrinsicData{ast::Intrinsic::kMin, ParamType::kU32, "metal::min"},
- IntrinsicData{ast::Intrinsic::kNormalize, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kLog, ParamType::kF32, "metal::log"},
+ IntrinsicData{semantic::Intrinsic::kLog2, ParamType::kF32,
+ "metal::log2"},
+ IntrinsicData{semantic::Intrinsic::kMax, ParamType::kF32,
+ "metal::fmax"},
+ IntrinsicData{semantic::Intrinsic::kMax, ParamType::kU32, "metal::max"},
+ 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{ast::Intrinsic::kReflect, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kPow, ParamType::kF32, "metal::pow"},
+ 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{ast::Intrinsic::kSelect, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kRound, ParamType::kU32,
+ "metal::round"},
+ IntrinsicData{semantic::Intrinsic::kSelect, ParamType::kF32,
"metal::select"},
- IntrinsicData{ast::Intrinsic::kSign, ParamType::kF32, "metal::sign"},
- IntrinsicData{ast::Intrinsic::kSin, ParamType::kF32, "metal::sin"},
- IntrinsicData{ast::Intrinsic::kSinh, ParamType::kF32, "metal::sinh"},
- IntrinsicData{ast::Intrinsic::kSmoothStep, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kSign, ParamType::kF32,
+ "metal::sign"},
+ IntrinsicData{semantic::Intrinsic::kSin, ParamType::kF32, "metal::sin"},
+ 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{ast::Intrinsic::kStep, ParamType::kF32, "metal::step"},
- IntrinsicData{ast::Intrinsic::kTan, ParamType::kF32, "metal::tan"},
- IntrinsicData{ast::Intrinsic::kTanh, ParamType::kF32, "metal::tanh"},
- IntrinsicData{ast::Intrinsic::kTrunc, ParamType::kF32,
+ IntrinsicData{semantic::Intrinsic::kSqrt, ParamType::kF32,
+ "metal::sqrt"},
+ IntrinsicData{semantic::Intrinsic::kStep, ParamType::kF32,
+ "metal::step"},
+ 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) {
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index 79c71e8..9374ac1 100644
--- a/src/writer/spirv/builder.cc
+++ b/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 @@
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 @@
} 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 @@
} 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 @@
} 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 @@
return 0;
}
- if (ident->IsIntrinsic()) {
- return GenerateIntrinsic(ident, expr);
+ auto* sem = builder_.Sem().Get(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::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 @@
return 0;
}
- auto intrinsic = ident->intrinsic();
+ auto intrinsic = sem->intrinsic();
- if (ast::intrinsic::IsFineDerivative(intrinsic) ||
- ast::intrinsic::IsCoarseDerivative(intrinsic)) {
+ if (semantic::intrinsic::IsFineDerivative(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 (!GenerateTextureIntrinsic(ident, call, Operand::Int(result_type_id),
- result)) {
+ if (auto* tex_sem = sem->As<semantic::TextureIntrinsicCall>()) {
+ if (!GenerateTextureIntrinsic(ident, call, tex_sem,
+ Operand::Int(result_type_id), result)) {
return 0;
}
return result_id;
@@ -1883,11 +1887,11 @@
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 @@
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 @@
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 @@
return result_id;
}
-bool Builder::GenerateTextureIntrinsic(ast::IdentifierExpression* ident,
- ast::CallExpression* call,
- Operand result_type,
- Operand result_id) {
- auto* sig = static_cast<const ast::intrinsic::TextureSignature*>(
- ident->intrinsic_signature());
- assert(sig != nullptr);
- auto& pidx = sig->params.idx;
- auto const kNotUsed = ast::intrinsic::TextureSignature::Parameters::kNotUsed;
+bool Builder::GenerateTextureIntrinsic(
+ ast::IdentifierExpression* ident,
+ ast::CallExpression* call,
+ const semantic::TextureIntrinsicCall* sem,
+ Operand result_type,
+ Operand result_id) {
+ auto& pidx = sem->Params().idx;
+ auto const kNotUsed = semantic::TextureIntrinsicCall::Parameters::kNotUsed;
assert(pidx.texture != kNotUsed);
auto* texture_type =
@@ -2149,8 +2152,8 @@
return append_coords_to_spirv_params();
};
- switch (ident->intrinsic()) {
- case ast::Intrinsic::kTextureDimensions: {
+ switch (sem->intrinsic()) {
+ 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 @@
}
break;
}
- case ast::Intrinsic::kTextureNumLayers: {
+ case semantic::Intrinsic::kTextureNumLayers: {
uint32_t spirv_dims = 0;
switch (texture_type->dim()) {
default:
@@ -2241,19 +2244,19 @@
}
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 @@
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 @@
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 @@
}
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 @@
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 @@
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 @@
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()) {
diff --git a/src/writer/spirv/builder.h b/src/writer/spirv/builder.h
index ee5f37f..3ab7479 100644
--- a/src/writer/spirv/builder.h
+++ b/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 @@
/// 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
