Rename semantic::Intrinsic to semantic::IntrinsicType
This allows us to create a semantic::Intrinsic class that holds more information about the particular intrinsic overload.
Change-Id: I180ddb507ebc92172badfdd3a59af346f96e1f02
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/40500
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/reader/spirv/function.cc b/src/reader/spirv/function.cc
index 0767da6..46fb806 100644
--- a/src/reader/spirv/function.cc
+++ b/src/reader/spirv/function.cc
@@ -470,19 +470,19 @@
}
// Returns the WGSL standard library function intrinsic for the
-// given instruction, or semantic::Intrinsic::kNone
-semantic::Intrinsic GetIntrinsic(SpvOp opcode) {
+// given instruction, or semantic::IntrinsicType::kNone
+semantic::IntrinsicType GetIntrinsic(SpvOp opcode) {
switch (opcode) {
case SpvOpBitCount:
- return semantic::Intrinsic::kCountOneBits;
+ return semantic::IntrinsicType::kCountOneBits;
case SpvOpBitReverse:
- return semantic::Intrinsic::kReverseBits;
+ return semantic::IntrinsicType::kReverseBits;
case SpvOpDot:
- return semantic::Intrinsic::kDot;
+ return semantic::IntrinsicType::kDot;
default:
break;
}
- return semantic::Intrinsic::kNone;
+ return semantic::IntrinsicType::kNone;
}
// @param opcode a SPIR-V opcode
@@ -3212,7 +3212,7 @@
}
const auto intrinsic = GetIntrinsic(opcode);
- if (intrinsic != semantic::Intrinsic::kNone) {
+ if (intrinsic != semantic::IntrinsicType::kNone) {
return MakeIntrinsicCall(inst);
}
diff --git a/src/semantic/call.h b/src/semantic/call.h
index c1459d5..6351421 100644
--- a/src/semantic/call.h
+++ b/src/semantic/call.h
@@ -40,16 +40,16 @@
/// Constructor
/// @param return_type the return type of the call
/// @param intrinsic the call target intrinsic
- IntrinsicCall(type::Type* return_type, Intrinsic intrinsic);
+ IntrinsicCall(type::Type* return_type, IntrinsicType intrinsic);
/// Destructor
~IntrinsicCall() override;
/// @returns the target intrinsic for the call
- Intrinsic intrinsic() const { return intrinsic_; }
+ IntrinsicType intrinsic() const { return intrinsic_; }
private:
- Intrinsic const intrinsic_;
+ IntrinsicType const intrinsic_;
};
/// TextureIntrinsicCall holds semantic information for ast::CallExpression
@@ -105,7 +105,7 @@
/// @param intrinsic the call target intrinsic
/// @param params the overload parameter info
TextureIntrinsicCall(type::Type* return_type,
- Intrinsic intrinsic,
+ IntrinsicType intrinsic,
const Parameters& params);
/// Destructor
diff --git a/src/semantic/intrinsic.cc b/src/semantic/intrinsic.cc
index 2ecd238..b2eea27 100644
--- a/src/semantic/intrinsic.cc
+++ b/src/semantic/intrinsic.cc
@@ -17,218 +17,221 @@
namespace tint {
namespace semantic {
-std::ostream& operator<<(std::ostream& out, Intrinsic i) {
+std::ostream& operator<<(std::ostream& out, IntrinsicType i) {
out << intrinsic::str(i);
return out;
}
namespace intrinsic {
-const char* str(Intrinsic i) {
+const char* str(IntrinsicType i) {
/// The emitted name matches the spelling in the WGSL spec.
/// including case.
switch (i) {
- case Intrinsic::kNone:
+ case IntrinsicType::kNone:
return "<not-an-intrinsic>";
- case Intrinsic::kAbs:
+ case IntrinsicType::kAbs:
return "abs";
- case Intrinsic::kAcos:
+ case IntrinsicType::kAcos:
return "acos";
- case Intrinsic::kAll:
+ case IntrinsicType::kAll:
return "all";
- case Intrinsic::kAny:
+ case IntrinsicType::kAny:
return "any";
- case Intrinsic::kArrayLength:
+ case IntrinsicType::kArrayLength:
return "arrayLength";
- case Intrinsic::kAsin:
+ case IntrinsicType::kAsin:
return "asin";
- case Intrinsic::kAtan:
+ case IntrinsicType::kAtan:
return "atan";
- case Intrinsic::kAtan2:
+ case IntrinsicType::kAtan2:
return "atan2";
- case Intrinsic::kCeil:
+ case IntrinsicType::kCeil:
return "ceil";
- case Intrinsic::kClamp:
+ case IntrinsicType::kClamp:
return "clamp";
- case Intrinsic::kCos:
+ case IntrinsicType::kCos:
return "cos";
- case Intrinsic::kCosh:
+ case IntrinsicType::kCosh:
return "cosh";
- case Intrinsic::kCountOneBits:
+ case IntrinsicType::kCountOneBits:
return "countOneBits";
- case Intrinsic::kCross:
+ case IntrinsicType::kCross:
return "cross";
- case Intrinsic::kDeterminant:
+ case IntrinsicType::kDeterminant:
return "determinant";
- case Intrinsic::kDistance:
+ case IntrinsicType::kDistance:
return "distance";
- case Intrinsic::kDot:
+ case IntrinsicType::kDot:
return "dot";
- case Intrinsic::kDpdx:
+ case IntrinsicType::kDpdx:
return "dpdx";
- case Intrinsic::kDpdxCoarse:
+ case IntrinsicType::kDpdxCoarse:
return "dpdxCoarse";
- case Intrinsic::kDpdxFine:
+ case IntrinsicType::kDpdxFine:
return "dpdxFine";
- case Intrinsic::kDpdy:
+ case IntrinsicType::kDpdy:
return "dpdy";
- case Intrinsic::kDpdyCoarse:
+ case IntrinsicType::kDpdyCoarse:
return "dpdyCoarse";
- case Intrinsic::kDpdyFine:
+ case IntrinsicType::kDpdyFine:
return "dpdyFine";
- case Intrinsic::kExp:
+ case IntrinsicType::kExp:
return "exp";
- case Intrinsic::kExp2:
+ case IntrinsicType::kExp2:
return "exp2";
- case Intrinsic::kFaceForward:
+ case IntrinsicType::kFaceForward:
return "faceForward";
- case Intrinsic::kFloor:
+ case IntrinsicType::kFloor:
return "floor";
- case Intrinsic::kFma:
+ case IntrinsicType::kFma:
return "fma";
- case Intrinsic::kFract:
+ case IntrinsicType::kFract:
return "fract";
- case Intrinsic::kFrexp:
+ case IntrinsicType::kFrexp:
return "frexp";
- case Intrinsic::kFwidth:
+ case IntrinsicType::kFwidth:
return "fwidth";
- case Intrinsic::kFwidthCoarse:
+ case IntrinsicType::kFwidthCoarse:
return "fwidthCoarse";
- case Intrinsic::kFwidthFine:
+ case IntrinsicType::kFwidthFine:
return "fwidthFine";
- case Intrinsic::kInverseSqrt:
+ case IntrinsicType::kInverseSqrt:
return "inverseSqrt";
- case Intrinsic::kIsFinite:
+ case IntrinsicType::kIsFinite:
return "isFinite";
- case Intrinsic::kIsInf:
+ case IntrinsicType::kIsInf:
return "isInf";
- case Intrinsic::kIsNan:
+ case IntrinsicType::kIsNan:
return "isNan";
- case Intrinsic::kIsNormal:
+ case IntrinsicType::kIsNormal:
return "isNormal";
- case Intrinsic::kLdexp:
+ case IntrinsicType::kLdexp:
return "ldexp";
- case Intrinsic::kLength:
+ case IntrinsicType::kLength:
return "length";
- case Intrinsic::kLog:
+ case IntrinsicType::kLog:
return "log";
- case Intrinsic::kLog2:
+ case IntrinsicType::kLog2:
return "log2";
- case Intrinsic::kMax:
+ case IntrinsicType::kMax:
return "max";
- case Intrinsic::kMin:
+ case IntrinsicType::kMin:
return "min";
- case Intrinsic::kMix:
+ case IntrinsicType::kMix:
return "mix";
- case Intrinsic::kModf:
+ case IntrinsicType::kModf:
return "modf";
- case Intrinsic::kNormalize:
+ case IntrinsicType::kNormalize:
return "normalize";
- case Intrinsic::kPack4x8Snorm:
+ case IntrinsicType::kPack4x8Snorm:
return "pack4x8snorm";
- case Intrinsic::kPack4x8Unorm:
+ case IntrinsicType::kPack4x8Unorm:
return "pack4x8unorm";
- case Intrinsic::kPack2x16Snorm:
+ case IntrinsicType::kPack2x16Snorm:
return "pack2x16snorm";
- case Intrinsic::kPack2x16Unorm:
+ case IntrinsicType::kPack2x16Unorm:
return "pack2x16unorm";
- case Intrinsic::kPack2x16Float:
+ case IntrinsicType::kPack2x16Float:
return "pack2x16float";
- case Intrinsic::kPow:
+ case IntrinsicType::kPow:
return "pow";
- case Intrinsic::kReflect:
+ case IntrinsicType::kReflect:
return "reflect";
- case Intrinsic::kReverseBits:
+ case IntrinsicType::kReverseBits:
return "reverseBits";
- case Intrinsic::kRound:
+ case IntrinsicType::kRound:
return "round";
- case Intrinsic::kSelect:
+ case IntrinsicType::kSelect:
return "select";
- case Intrinsic::kSign:
+ case IntrinsicType::kSign:
return "sign";
- case Intrinsic::kSin:
+ case IntrinsicType::kSin:
return "sin";
- case Intrinsic::kSinh:
+ case IntrinsicType::kSinh:
return "sinh";
- case Intrinsic::kSmoothStep:
+ case IntrinsicType::kSmoothStep:
return "smoothStep";
- case Intrinsic::kSqrt:
+ case IntrinsicType::kSqrt:
return "sqrt";
- case Intrinsic::kStep:
+ case IntrinsicType::kStep:
return "step";
- case Intrinsic::kTan:
+ case IntrinsicType::kTan:
return "tan";
- case Intrinsic::kTanh:
+ case IntrinsicType::kTanh:
return "tanh";
- case Intrinsic::kTextureDimensions:
+ case IntrinsicType::kTextureDimensions:
return "textureDimensions";
- case Intrinsic::kTextureLoad:
+ case IntrinsicType::kTextureLoad:
return "textureLoad";
- case Intrinsic::kTextureNumLayers:
+ case IntrinsicType::kTextureNumLayers:
return "textureNumLayers";
- case Intrinsic::kTextureNumLevels:
+ case IntrinsicType::kTextureNumLevels:
return "textureNumLevels";
- case Intrinsic::kTextureNumSamples:
+ case IntrinsicType::kTextureNumSamples:
return "textureNumSamples";
- case Intrinsic::kTextureSample:
+ case IntrinsicType::kTextureSample:
return "textureSample";
- case Intrinsic::kTextureSampleBias:
+ case IntrinsicType::kTextureSampleBias:
return "textureSampleBias";
- case Intrinsic::kTextureSampleCompare:
+ case IntrinsicType::kTextureSampleCompare:
return "textureSampleCompare";
- case Intrinsic::kTextureSampleGrad:
+ case IntrinsicType::kTextureSampleGrad:
return "textureSampleGrad";
- case Intrinsic::kTextureSampleLevel:
+ case IntrinsicType::kTextureSampleLevel:
return "textureSampleLevel";
- case Intrinsic::kTextureStore:
+ case IntrinsicType::kTextureStore:
return "textureStore";
- case Intrinsic::kTrunc:
+ case IntrinsicType::kTrunc:
return "trunc";
}
return "<unknown>";
}
-bool IsCoarseDerivative(Intrinsic i) {
- return i == Intrinsic::kDpdxCoarse || i == Intrinsic::kDpdyCoarse ||
- i == Intrinsic::kFwidthCoarse;
+bool IsCoarseDerivative(IntrinsicType i) {
+ return i == IntrinsicType::kDpdxCoarse || i == IntrinsicType::kDpdyCoarse ||
+ i == IntrinsicType::kFwidthCoarse;
}
-bool IsFineDerivative(Intrinsic i) {
- return i == Intrinsic::kDpdxFine || i == Intrinsic::kDpdyFine ||
- i == Intrinsic::kFwidthFine;
+bool IsFineDerivative(IntrinsicType i) {
+ return i == IntrinsicType::kDpdxFine || i == IntrinsicType::kDpdyFine ||
+ i == IntrinsicType::kFwidthFine;
}
-bool IsDerivative(Intrinsic i) {
- return i == Intrinsic::kDpdx || i == Intrinsic::kDpdy ||
- i == Intrinsic::kFwidth || IsCoarseDerivative(i) ||
+bool IsDerivative(IntrinsicType i) {
+ return i == IntrinsicType::kDpdx || i == IntrinsicType::kDpdy ||
+ i == IntrinsicType::kFwidth || IsCoarseDerivative(i) ||
IsFineDerivative(i);
}
-bool IsFloatClassificationIntrinsic(Intrinsic i) {
- return i == Intrinsic::kIsFinite || i == Intrinsic::kIsInf ||
- i == Intrinsic::kIsNan || i == Intrinsic::kIsNormal;
+bool IsFloatClassificationIntrinsic(IntrinsicType i) {
+ return i == IntrinsicType::kIsFinite || i == IntrinsicType::kIsInf ||
+ i == IntrinsicType::kIsNan || i == IntrinsicType::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 IsTextureIntrinsic(IntrinsicType i) {
+ return IsImageQueryIntrinsic(i) || i == IntrinsicType::kTextureLoad ||
+ i == IntrinsicType::kTextureSample ||
+ i == IntrinsicType::kTextureSampleLevel ||
+ i == IntrinsicType::kTextureSampleBias ||
+ i == IntrinsicType::kTextureSampleCompare ||
+ i == IntrinsicType::kTextureSampleGrad ||
+ i == IntrinsicType::kTextureStore;
}
-bool IsImageQueryIntrinsic(Intrinsic i) {
- return i == semantic::Intrinsic::kTextureDimensions ||
- i == Intrinsic::kTextureNumLayers ||
- i == Intrinsic::kTextureNumLevels ||
- i == Intrinsic::kTextureNumSamples;
+bool IsImageQueryIntrinsic(IntrinsicType i) {
+ return i == semantic::IntrinsicType::kTextureDimensions ||
+ i == IntrinsicType::kTextureNumLayers ||
+ i == IntrinsicType::kTextureNumLevels ||
+ i == IntrinsicType::kTextureNumSamples;
}
-bool IsDataPackingIntrinsic(Intrinsic i) {
- return i == Intrinsic::kPack4x8Snorm || i == Intrinsic::kPack4x8Unorm ||
- i == Intrinsic::kPack2x16Snorm || i == Intrinsic::kPack2x16Unorm ||
- i == Intrinsic::kPack2x16Float;
+bool IsDataPackingIntrinsic(IntrinsicType i) {
+ return i == IntrinsicType::kPack4x8Snorm ||
+ i == IntrinsicType::kPack4x8Unorm ||
+ i == IntrinsicType::kPack2x16Snorm ||
+ i == IntrinsicType::kPack2x16Unorm ||
+ i == IntrinsicType::kPack2x16Float;
}
} // namespace intrinsic
diff --git a/src/semantic/intrinsic.h b/src/semantic/intrinsic.h
index fd5a544..f006f31 100644
--- a/src/semantic/intrinsic.h
+++ b/src/semantic/intrinsic.h
@@ -20,7 +20,7 @@
namespace tint {
namespace semantic {
-enum class Intrinsic {
+enum class IntrinsicType {
kNone = -1,
kAbs,
@@ -104,48 +104,48 @@
/// 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);
+std::ostream& operator<<(std::ostream& out, IntrinsicType 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);
+bool IsCoarseDerivative(IntrinsicType 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);
+bool IsFineDerivative(IntrinsicType 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);
+bool IsDerivative(IntrinsicType 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);
+bool IsFloatClassificationIntrinsic(IntrinsicType 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);
+bool IsTextureIntrinsic(IntrinsicType 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);
+bool IsImageQueryIntrinsic(IntrinsicType i);
/// Determines if the given `i` is a data packing intrinsic
/// @param i the intrinsic
/// @returns true if the given `i` is a data packing intrinsic
-bool IsDataPackingIntrinsic(Intrinsic i);
+bool IsDataPackingIntrinsic(IntrinsicType i);
/// @returns the name of the intrinsic function. The spelling, including case,
/// matches the name in the WGSL spec.
-const char* str(Intrinsic i);
+const char* str(IntrinsicType i);
} // namespace intrinsic
} // namespace semantic
diff --git a/src/semantic/sem_call.cc b/src/semantic/sem_call.cc
index 72be60a..777ad51 100644
--- a/src/semantic/sem_call.cc
+++ b/src/semantic/sem_call.cc
@@ -25,13 +25,13 @@
Call::~Call() = default;
-IntrinsicCall::IntrinsicCall(type::Type* return_type, Intrinsic intrinsic)
+IntrinsicCall::IntrinsicCall(type::Type* return_type, IntrinsicType intrinsic)
: Base(return_type), intrinsic_(intrinsic) {}
IntrinsicCall::~IntrinsicCall() = default;
TextureIntrinsicCall::TextureIntrinsicCall(type::Type* return_type,
- Intrinsic intrinsic,
+ IntrinsicType intrinsic,
const Parameters& params)
: Base(return_type, intrinsic), params_(params) {}
diff --git a/src/type_determiner.cc b/src/type_determiner.cc
index 2e72f4a..6b47ffa 100644
--- a/src/type_determiner.cc
+++ b/src/type_determiner.cc
@@ -65,6 +65,11 @@
#include "src/type/void_type.h"
namespace tint {
+namespace {
+
+using IntrinsicType = tint::semantic::IntrinsicType;
+
+} // namespace
TypeDeterminer::TypeDeterminer(ProgramBuilder* builder) : builder_(builder) {}
@@ -415,7 +420,7 @@
auto name = builder_->Symbols().NameFor(ident->symbol());
auto intrinsic = MatchIntrinsic(name);
- if (intrinsic != semantic::Intrinsic::kNone) {
+ if (intrinsic != IntrinsicType::kNone) {
if (!DetermineIntrinsicCall(call, intrinsic)) {
return false;
}
@@ -464,7 +469,7 @@
};
struct IntrinsicData {
- semantic::Intrinsic intrinsic;
+ IntrinsicType intrinsic;
IntrinsicDataType result_type;
uint8_t result_vector_width;
uint8_t param_for_result_type;
@@ -473,74 +478,68 @@
// 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[] = {
- {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::kPack4x8Snorm, IntrinsicDataType::kUnsignedInteger, 1,
- 0},
- {semantic::Intrinsic::kPack4x8Unorm, IntrinsicDataType::kUnsignedInteger, 1,
- 0},
- {semantic::Intrinsic::kPack2x16Snorm, IntrinsicDataType::kUnsignedInteger,
- 1, 0},
- {semantic::Intrinsic::kPack2x16Unorm, IntrinsicDataType::kUnsignedInteger,
- 1, 0},
- {semantic::Intrinsic::kPack2x16Float, IntrinsicDataType::kUnsignedInteger,
- 1, 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},
+ {IntrinsicType::kAbs, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kAcos, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kAll, IntrinsicDataType::kBool, 1, 0},
+ {IntrinsicType::kAny, IntrinsicDataType::kBool, 1, 0},
+ {IntrinsicType::kArrayLength, IntrinsicDataType::kUnsignedInteger, 1, 0},
+ {IntrinsicType::kAsin, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kAtan, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kAtan2, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kCeil, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kClamp, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kCos, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kCosh, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kCountOneBits, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kCross, IntrinsicDataType::kFloat, 3, 0},
+ {IntrinsicType::kDeterminant, IntrinsicDataType::kFloat, 1, 0},
+ {IntrinsicType::kDistance, IntrinsicDataType::kFloat, 1, 0},
+ {IntrinsicType::kDpdx, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kDpdxCoarse, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kDpdxFine, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kDpdy, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kDpdyCoarse, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kDpdyFine, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kDot, IntrinsicDataType::kFloat, 1, 0},
+ {IntrinsicType::kExp, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kExp2, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kFaceForward, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kFloor, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kFwidth, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kFwidthCoarse, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kFwidthFine, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kFma, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kFract, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kFrexp, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kInverseSqrt, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kLdexp, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kLength, IntrinsicDataType::kFloat, 1, 0},
+ {IntrinsicType::kLog, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kLog2, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kMax, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kMin, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kMix, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kModf, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kNormalize, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kPack4x8Snorm, IntrinsicDataType::kUnsignedInteger, 1, 0},
+ {IntrinsicType::kPack4x8Unorm, IntrinsicDataType::kUnsignedInteger, 1, 0},
+ {IntrinsicType::kPack2x16Snorm, IntrinsicDataType::kUnsignedInteger, 1, 0},
+ {IntrinsicType::kPack2x16Unorm, IntrinsicDataType::kUnsignedInteger, 1, 0},
+ {IntrinsicType::kPack2x16Float, IntrinsicDataType::kUnsignedInteger, 1, 0},
+ {IntrinsicType::kPow, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kReflect, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kReverseBits, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kRound, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kSelect, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kSign, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kSin, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kSinh, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kSmoothStep, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kSqrt, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kStep, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kTan, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kTanh, IntrinsicDataType::kDependent, 0, 0},
+ {IntrinsicType::kTrunc, IntrinsicDataType::kDependent, 0, 0},
};
constexpr const uint32_t kIntrinsicDataCount =
@@ -549,7 +548,7 @@
} // namespace
bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
- semantic::Intrinsic intrinsic) {
+ IntrinsicType intrinsic) {
auto create_sem = [&](type::Type* result) {
auto* sem = builder_->create<semantic::IntrinsicCall>(result, intrinsic);
builder_->Sem().Add(call, sem);
@@ -586,18 +585,18 @@
bool is_array = type::IsTextureArray(texture->dim());
bool is_multisampled = texture->Is<type::MultisampledTexture>();
switch (intrinsic) {
- case semantic::Intrinsic::kTextureDimensions:
+ case IntrinsicType::kTextureDimensions:
param.idx.texture = param.count++;
if (call->params().size() > param.count) {
param.idx.level = param.count++;
}
break;
- case semantic::Intrinsic::kTextureNumLayers:
- case semantic::Intrinsic::kTextureNumLevels:
- case semantic::Intrinsic::kTextureNumSamples:
+ case IntrinsicType::kTextureNumLayers:
+ case IntrinsicType::kTextureNumLevels:
+ case IntrinsicType::kTextureNumSamples:
param.idx.texture = param.count++;
break;
- case semantic::Intrinsic::kTextureLoad:
+ case IntrinsicType::kTextureLoad:
param.idx.texture = param.count++;
param.idx.coords = param.count++;
if (is_array) {
@@ -611,7 +610,7 @@
}
}
break;
- case semantic::Intrinsic::kTextureSample:
+ case IntrinsicType::kTextureSample:
param.idx.texture = param.count++;
param.idx.sampler = param.count++;
param.idx.coords = param.count++;
@@ -622,7 +621,7 @@
param.idx.offset = param.count++;
}
break;
- case semantic::Intrinsic::kTextureSampleBias:
+ case IntrinsicType::kTextureSampleBias:
param.idx.texture = param.count++;
param.idx.sampler = param.count++;
param.idx.coords = param.count++;
@@ -634,7 +633,7 @@
param.idx.offset = param.count++;
}
break;
- case semantic::Intrinsic::kTextureSampleLevel:
+ case IntrinsicType::kTextureSampleLevel:
param.idx.texture = param.count++;
param.idx.sampler = param.count++;
param.idx.coords = param.count++;
@@ -646,7 +645,7 @@
param.idx.offset = param.count++;
}
break;
- case semantic::Intrinsic::kTextureSampleCompare:
+ case IntrinsicType::kTextureSampleCompare:
param.idx.texture = param.count++;
param.idx.sampler = param.count++;
param.idx.coords = param.count++;
@@ -658,7 +657,7 @@
param.idx.offset = param.count++;
}
break;
- case semantic::Intrinsic::kTextureSampleGrad:
+ case IntrinsicType::kTextureSampleGrad:
param.idx.texture = param.count++;
param.idx.sampler = param.count++;
param.idx.coords = param.count++;
@@ -671,7 +670,7 @@
param.idx.offset = param.count++;
}
break;
- case semantic::Intrinsic::kTextureStore:
+ case IntrinsicType::kTextureStore:
param.idx.texture = param.count++;
param.idx.coords = param.count++;
if (is_array) {
@@ -697,7 +696,7 @@
// Set the function return type
type::Type* return_type = nullptr;
switch (intrinsic) {
- case semantic::Intrinsic::kTextureDimensions: {
+ case IntrinsicType::kTextureDimensions: {
auto* i32 = builder_->create<type::I32>();
switch (texture->dim()) {
default:
@@ -719,12 +718,12 @@
}
break;
}
- case semantic::Intrinsic::kTextureNumLayers:
- case semantic::Intrinsic::kTextureNumLevels:
- case semantic::Intrinsic::kTextureNumSamples:
+ case IntrinsicType::kTextureNumLayers:
+ case IntrinsicType::kTextureNumLevels:
+ case IntrinsicType::kTextureNumSamples:
return_type = builder_->create<type::I32>();
break;
- case semantic::Intrinsic::kTextureStore:
+ case IntrinsicType::kTextureStore:
return_type = builder_->create<type::Void>();
break;
default: {
@@ -848,7 +847,7 @@
}
std::string name = builder_->Symbols().NameFor(symbol);
- if (MatchIntrinsic(name) != semantic::Intrinsic::kNone) {
+ if (MatchIntrinsic(name) != IntrinsicType::kNone) {
// Identifier is to an intrinsic function, which has no type (currently).
return true;
}
@@ -858,163 +857,163 @@
return false;
}
-semantic::Intrinsic TypeDeterminer::MatchIntrinsic(const std::string& name) {
+IntrinsicType TypeDeterminer::MatchIntrinsic(const std::string& name) {
if (name == "abs") {
- return semantic::Intrinsic::kAbs;
+ return IntrinsicType::kAbs;
} else if (name == "acos") {
- return semantic::Intrinsic::kAcos;
+ return IntrinsicType::kAcos;
} else if (name == "all") {
- return semantic::Intrinsic::kAll;
+ return IntrinsicType::kAll;
} else if (name == "any") {
- return semantic::Intrinsic::kAny;
+ return IntrinsicType::kAny;
} else if (name == "arrayLength") {
- return semantic::Intrinsic::kArrayLength;
+ return IntrinsicType::kArrayLength;
} else if (name == "asin") {
- return semantic::Intrinsic::kAsin;
+ return IntrinsicType::kAsin;
} else if (name == "atan") {
- return semantic::Intrinsic::kAtan;
+ return IntrinsicType::kAtan;
} else if (name == "atan2") {
- return semantic::Intrinsic::kAtan2;
+ return IntrinsicType::kAtan2;
} else if (name == "ceil") {
- return semantic::Intrinsic::kCeil;
+ return IntrinsicType::kCeil;
} else if (name == "clamp") {
- return semantic::Intrinsic::kClamp;
+ return IntrinsicType::kClamp;
} else if (name == "cos") {
- return semantic::Intrinsic::kCos;
+ return IntrinsicType::kCos;
} else if (name == "cosh") {
- return semantic::Intrinsic::kCosh;
+ return IntrinsicType::kCosh;
} else if (name == "countOneBits") {
- return semantic::Intrinsic::kCountOneBits;
+ return IntrinsicType::kCountOneBits;
} else if (name == "cross") {
- return semantic::Intrinsic::kCross;
+ return IntrinsicType::kCross;
} else if (name == "determinant") {
- return semantic::Intrinsic::kDeterminant;
+ return IntrinsicType::kDeterminant;
} else if (name == "distance") {
- return semantic::Intrinsic::kDistance;
+ return IntrinsicType::kDistance;
} else if (name == "dot") {
- return semantic::Intrinsic::kDot;
+ return IntrinsicType::kDot;
} else if (name == "dpdx") {
- return semantic::Intrinsic::kDpdx;
+ return IntrinsicType::kDpdx;
} else if (name == "dpdxCoarse") {
- return semantic::Intrinsic::kDpdxCoarse;
+ return IntrinsicType::kDpdxCoarse;
} else if (name == "dpdxFine") {
- return semantic::Intrinsic::kDpdxFine;
+ return IntrinsicType::kDpdxFine;
} else if (name == "dpdy") {
- return semantic::Intrinsic::kDpdy;
+ return IntrinsicType::kDpdy;
} else if (name == "dpdyCoarse") {
- return semantic::Intrinsic::kDpdyCoarse;
+ return IntrinsicType::kDpdyCoarse;
} else if (name == "dpdyFine") {
- return semantic::Intrinsic::kDpdyFine;
+ return IntrinsicType::kDpdyFine;
} else if (name == "exp") {
- return semantic::Intrinsic::kExp;
+ return IntrinsicType::kExp;
} else if (name == "exp2") {
- return semantic::Intrinsic::kExp2;
+ return IntrinsicType::kExp2;
} else if (name == "faceForward") {
- return semantic::Intrinsic::kFaceForward;
+ return IntrinsicType::kFaceForward;
} else if (name == "floor") {
- return semantic::Intrinsic::kFloor;
+ return IntrinsicType::kFloor;
} else if (name == "fma") {
- return semantic::Intrinsic::kFma;
+ return IntrinsicType::kFma;
} else if (name == "fract") {
- return semantic::Intrinsic::kFract;
+ return IntrinsicType::kFract;
} else if (name == "frexp") {
- return semantic::Intrinsic::kFrexp;
+ return IntrinsicType::kFrexp;
} else if (name == "fwidth") {
- return semantic::Intrinsic::kFwidth;
+ return IntrinsicType::kFwidth;
} else if (name == "fwidthCoarse") {
- return semantic::Intrinsic::kFwidthCoarse;
+ return IntrinsicType::kFwidthCoarse;
} else if (name == "fwidthFine") {
- return semantic::Intrinsic::kFwidthFine;
+ return IntrinsicType::kFwidthFine;
} else if (name == "inverseSqrt") {
- return semantic::Intrinsic::kInverseSqrt;
+ return IntrinsicType::kInverseSqrt;
} else if (name == "isFinite") {
- return semantic::Intrinsic::kIsFinite;
+ return IntrinsicType::kIsFinite;
} else if (name == "isInf") {
- return semantic::Intrinsic::kIsInf;
+ return IntrinsicType::kIsInf;
} else if (name == "isNan") {
- return semantic::Intrinsic::kIsNan;
+ return IntrinsicType::kIsNan;
} else if (name == "isNormal") {
- return semantic::Intrinsic::kIsNormal;
+ return IntrinsicType::kIsNormal;
} else if (name == "ldexp") {
- return semantic::Intrinsic::kLdexp;
+ return IntrinsicType::kLdexp;
} else if (name == "length") {
- return semantic::Intrinsic::kLength;
+ return IntrinsicType::kLength;
} else if (name == "log") {
- return semantic::Intrinsic::kLog;
+ return IntrinsicType::kLog;
} else if (name == "log2") {
- return semantic::Intrinsic::kLog2;
+ return IntrinsicType::kLog2;
} else if (name == "max") {
- return semantic::Intrinsic::kMax;
+ return IntrinsicType::kMax;
} else if (name == "min") {
- return semantic::Intrinsic::kMin;
+ return IntrinsicType::kMin;
} else if (name == "mix") {
- return semantic::Intrinsic::kMix;
+ return IntrinsicType::kMix;
} else if (name == "modf") {
- return semantic::Intrinsic::kModf;
+ return IntrinsicType::kModf;
} else if (name == "normalize") {
- return semantic::Intrinsic::kNormalize;
+ return IntrinsicType::kNormalize;
} else if (name == "pack4x8snorm") {
- return semantic::Intrinsic::kPack4x8Snorm;
+ return IntrinsicType::kPack4x8Snorm;
} else if (name == "pack4x8unorm") {
- return semantic::Intrinsic::kPack4x8Unorm;
+ return IntrinsicType::kPack4x8Unorm;
} else if (name == "pack2x16snorm") {
- return semantic::Intrinsic::kPack2x16Snorm;
+ return IntrinsicType::kPack2x16Snorm;
} else if (name == "pack2x16unorm") {
- return semantic::Intrinsic::kPack2x16Unorm;
+ return IntrinsicType::kPack2x16Unorm;
} else if (name == "pack2x16float") {
- return semantic::Intrinsic::kPack2x16Float;
+ return IntrinsicType::kPack2x16Float;
} else if (name == "pow") {
- return semantic::Intrinsic::kPow;
+ return IntrinsicType::kPow;
} else if (name == "reflect") {
- return semantic::Intrinsic::kReflect;
+ return IntrinsicType::kReflect;
} else if (name == "reverseBits") {
- return semantic::Intrinsic::kReverseBits;
+ return IntrinsicType::kReverseBits;
} else if (name == "round") {
- return semantic::Intrinsic::kRound;
+ return IntrinsicType::kRound;
} else if (name == "select") {
- return semantic::Intrinsic::kSelect;
+ return IntrinsicType::kSelect;
} else if (name == "sign") {
- return semantic::Intrinsic::kSign;
+ return IntrinsicType::kSign;
} else if (name == "sin") {
- return semantic::Intrinsic::kSin;
+ return IntrinsicType::kSin;
} else if (name == "sinh") {
- return semantic::Intrinsic::kSinh;
+ return IntrinsicType::kSinh;
} else if (name == "smoothStep") {
- return semantic::Intrinsic::kSmoothStep;
+ return IntrinsicType::kSmoothStep;
} else if (name == "sqrt") {
- return semantic::Intrinsic::kSqrt;
+ return IntrinsicType::kSqrt;
} else if (name == "step") {
- return semantic::Intrinsic::kStep;
+ return IntrinsicType::kStep;
} else if (name == "tan") {
- return semantic::Intrinsic::kTan;
+ return IntrinsicType::kTan;
} else if (name == "tanh") {
- return semantic::Intrinsic::kTanh;
+ return IntrinsicType::kTanh;
} else if (name == "textureDimensions") {
- return semantic::Intrinsic::kTextureDimensions;
+ return IntrinsicType::kTextureDimensions;
} else if (name == "textureNumLayers") {
- return semantic::Intrinsic::kTextureNumLayers;
+ return IntrinsicType::kTextureNumLayers;
} else if (name == "textureNumLevels") {
- return semantic::Intrinsic::kTextureNumLevels;
+ return IntrinsicType::kTextureNumLevels;
} else if (name == "textureNumSamples") {
- return semantic::Intrinsic::kTextureNumSamples;
+ return IntrinsicType::kTextureNumSamples;
} else if (name == "textureLoad") {
- return semantic::Intrinsic::kTextureLoad;
+ return IntrinsicType::kTextureLoad;
} else if (name == "textureStore") {
- return semantic::Intrinsic::kTextureStore;
+ return IntrinsicType::kTextureStore;
} else if (name == "textureSample") {
- return semantic::Intrinsic::kTextureSample;
+ return IntrinsicType::kTextureSample;
} else if (name == "textureSampleBias") {
- return semantic::Intrinsic::kTextureSampleBias;
+ return IntrinsicType::kTextureSampleBias;
} else if (name == "textureSampleCompare") {
- return semantic::Intrinsic::kTextureSampleCompare;
+ return IntrinsicType::kTextureSampleCompare;
} else if (name == "textureSampleGrad") {
- return semantic::Intrinsic::kTextureSampleGrad;
+ return IntrinsicType::kTextureSampleGrad;
} else if (name == "textureSampleLevel") {
- return semantic::Intrinsic::kTextureSampleLevel;
+ return IntrinsicType::kTextureSampleLevel;
} else if (name == "trunc") {
- return semantic::Intrinsic::kTrunc;
+ return IntrinsicType::kTrunc;
}
- return semantic::Intrinsic::kNone;
+ return IntrinsicType::kNone;
}
bool TypeDeterminer::DetermineMemberAccessor(
diff --git a/src/type_determiner.h b/src/type_determiner.h
index 916a6b2..c8dc34d 100644
--- a/src/type_determiner.h
+++ b/src/type_determiner.h
@@ -67,9 +67,9 @@
bool Determine();
/// @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);
+ /// @return the semantic::IntrinsicType for the given name. If `name` does not
+ /// match an intrinsic, returns semantic::IntrinsicType::kNone
+ static semantic::IntrinsicType MatchIntrinsic(const std::string& name);
private:
template <typename T>
@@ -177,7 +177,7 @@
bool DetermineConstructor(ast::ConstructorExpression* expr);
bool DetermineIdentifier(ast::IdentifierExpression* expr);
bool DetermineIntrinsicCall(ast::CallExpression* call,
- semantic::Intrinsic intrinsic);
+ semantic::IntrinsicType 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 c6822c1..9013905 100644
--- a/src/type_determiner_test.cc
+++ b/src/type_determiner_test.cc
@@ -75,6 +75,8 @@
namespace tint {
namespace {
+using IntrinsicType = semantic::IntrinsicType;
+
class FakeStmt : public ast::Statement {
public:
explicit FakeStmt(Source source) : ast::Statement(source) {}
@@ -1618,7 +1620,7 @@
struct IntrinsicData {
const char* name;
- semantic::Intrinsic intrinsic;
+ IntrinsicType intrinsic;
};
inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
out << data.name;
@@ -1634,89 +1636,82 @@
TypeDeterminerTest,
IntrinsicDataTest,
testing::Values(
- 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{"abs", IntrinsicType::kAbs},
+ IntrinsicData{"acos", IntrinsicType::kAcos},
+ IntrinsicData{"all", IntrinsicType::kAll},
+ IntrinsicData{"any", IntrinsicType::kAny},
+ IntrinsicData{"arrayLength", IntrinsicType::kArrayLength},
+ IntrinsicData{"asin", IntrinsicType::kAsin},
+ IntrinsicData{"atan", IntrinsicType::kAtan},
+ IntrinsicData{"atan2", IntrinsicType::kAtan2},
+ IntrinsicData{"ceil", IntrinsicType::kCeil},
+ IntrinsicData{"clamp", IntrinsicType::kClamp},
+ IntrinsicData{"cos", IntrinsicType::kCos},
+ IntrinsicData{"cosh", IntrinsicType::kCosh},
+ IntrinsicData{"countOneBits", IntrinsicType::kCountOneBits},
+ IntrinsicData{"cross", IntrinsicType::kCross},
+ IntrinsicData{"determinant", IntrinsicType::kDeterminant},
+ IntrinsicData{"distance", IntrinsicType::kDistance},
+ IntrinsicData{"dot", IntrinsicType::kDot},
+ IntrinsicData{"dpdx", IntrinsicType::kDpdx},
+ IntrinsicData{"dpdxCoarse", IntrinsicType::kDpdxCoarse},
+ IntrinsicData{"dpdxFine", IntrinsicType::kDpdxFine},
+ IntrinsicData{"dpdy", IntrinsicType::kDpdy},
+ IntrinsicData{"dpdyCoarse", IntrinsicType::kDpdyCoarse},
+ IntrinsicData{"dpdyFine", IntrinsicType::kDpdyFine},
+ IntrinsicData{"exp", IntrinsicType::kExp},
+ IntrinsicData{"exp2", IntrinsicType::kExp2},
+ IntrinsicData{"faceForward", IntrinsicType::kFaceForward},
+ IntrinsicData{"floor", IntrinsicType::kFloor},
+ IntrinsicData{"fma", IntrinsicType::kFma},
+ IntrinsicData{"fract", IntrinsicType::kFract},
+ IntrinsicData{"frexp", IntrinsicType::kFrexp},
+ IntrinsicData{"fwidth", IntrinsicType::kFwidth},
+ IntrinsicData{"fwidthCoarse", IntrinsicType::kFwidthCoarse},
+ IntrinsicData{"fwidthFine", IntrinsicType::kFwidthFine},
+ IntrinsicData{"inverseSqrt", IntrinsicType::kInverseSqrt},
+ IntrinsicData{"isFinite", IntrinsicType::kIsFinite},
+ IntrinsicData{"isInf", IntrinsicType::kIsInf},
+ IntrinsicData{"isNan", IntrinsicType::kIsNan},
+ IntrinsicData{"isNormal", IntrinsicType::kIsNormal},
+ IntrinsicData{"ldexp", IntrinsicType::kLdexp},
+ IntrinsicData{"length", IntrinsicType::kLength},
+ IntrinsicData{"log", IntrinsicType::kLog},
+ IntrinsicData{"log2", IntrinsicType::kLog2},
+ IntrinsicData{"max", IntrinsicType::kMax},
+ IntrinsicData{"min", IntrinsicType::kMin},
+ IntrinsicData{"mix", IntrinsicType::kMix},
+ IntrinsicData{"modf", IntrinsicType::kModf},
+ IntrinsicData{"normalize", IntrinsicType::kNormalize},
+ IntrinsicData{"pow", IntrinsicType::kPow},
+ IntrinsicData{"reflect", IntrinsicType::kReflect},
+ IntrinsicData{"reverseBits", IntrinsicType::kReverseBits},
+ IntrinsicData{"round", IntrinsicType::kRound},
+ IntrinsicData{"select", IntrinsicType::kSelect},
+ IntrinsicData{"sign", IntrinsicType::kSign},
+ IntrinsicData{"sin", IntrinsicType::kSin},
+ IntrinsicData{"sinh", IntrinsicType::kSinh},
+ IntrinsicData{"smoothStep", IntrinsicType::kSmoothStep},
+ IntrinsicData{"sqrt", IntrinsicType::kSqrt},
+ IntrinsicData{"step", IntrinsicType::kStep},
+ IntrinsicData{"tan", IntrinsicType::kTan},
+ IntrinsicData{"tanh", IntrinsicType::kTanh},
+ IntrinsicData{"textureDimensions", IntrinsicType::kTextureDimensions},
+ IntrinsicData{"textureLoad", IntrinsicType::kTextureLoad},
+ IntrinsicData{"textureNumLayers", IntrinsicType::kTextureNumLayers},
+ IntrinsicData{"textureNumLevels", IntrinsicType::kTextureNumLevels},
+ IntrinsicData{"textureNumSamples", IntrinsicType::kTextureNumSamples},
+ IntrinsicData{"textureSample", IntrinsicType::kTextureSample},
+ IntrinsicData{"textureSampleBias", IntrinsicType::kTextureSampleBias},
IntrinsicData{"textureSampleCompare",
- semantic::Intrinsic::kTextureSampleCompare},
- IntrinsicData{"textureSampleGrad",
- semantic::Intrinsic::kTextureSampleGrad},
- IntrinsicData{"textureSampleLevel",
- semantic::Intrinsic::kTextureSampleLevel},
- IntrinsicData{"trunc", semantic::Intrinsic::kTrunc}));
+ IntrinsicType::kTextureSampleCompare},
+ IntrinsicData{"textureSampleGrad", IntrinsicType::kTextureSampleGrad},
+ IntrinsicData{"textureSampleLevel", IntrinsicType::kTextureSampleLevel},
+ IntrinsicData{"trunc", IntrinsicType::kTrunc}));
TEST_F(TypeDeterminerTest, MatchIntrinsicNoMatch) {
EXPECT_EQ(TypeDeterminer::MatchIntrinsic("not_intrinsic"),
- semantic::Intrinsic::kNone);
+ IntrinsicType::kNone);
}
using ImportData_DataPackingTest = TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1736,11 +1731,11 @@
TypeDeterminerTest,
ImportData_DataPackingTest,
testing::Values(
- IntrinsicData{"pack4x8snorm", semantic::Intrinsic::kPack4x8Snorm},
- IntrinsicData{"pack4x8unorm", semantic::Intrinsic::kPack4x8Unorm},
- IntrinsicData{"pack2x16snorm", semantic::Intrinsic::kPack2x16Snorm},
- IntrinsicData{"pack2x16unorm", semantic::Intrinsic::kPack2x16Unorm},
- IntrinsicData{"pack2x16float", semantic::Intrinsic::kPack2x16Float}));
+ IntrinsicData{"pack4x8snorm", IntrinsicType::kPack4x8Snorm},
+ IntrinsicData{"pack4x8unorm", IntrinsicType::kPack4x8Unorm},
+ IntrinsicData{"pack2x16snorm", IntrinsicType::kPack2x16Snorm},
+ IntrinsicData{"pack2x16unorm", IntrinsicType::kPack2x16Unorm},
+ IntrinsicData{"pack2x16float", IntrinsicType::kPack2x16Float}));
using ImportData_SingleParamTest = TypeDeterminerTestWithParam<IntrinsicData>;
TEST_P(ImportData_SingleParamTest, Scalar) {
@@ -1786,29 +1781,28 @@
INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest,
ImportData_SingleParamTest,
- 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}));
+ testing::Values(IntrinsicData{"acos", IntrinsicType::kAcos},
+ IntrinsicData{"asin", IntrinsicType::kAsin},
+ IntrinsicData{"atan", IntrinsicType::kAtan},
+ IntrinsicData{"ceil", IntrinsicType::kCeil},
+ IntrinsicData{"cos", IntrinsicType::kCos},
+ IntrinsicData{"cosh", IntrinsicType::kCosh},
+ IntrinsicData{"exp", IntrinsicType::kExp},
+ IntrinsicData{"exp2", IntrinsicType::kExp2},
+ IntrinsicData{"floor", IntrinsicType::kFloor},
+ IntrinsicData{"fract", IntrinsicType::kFract},
+ IntrinsicData{"inverseSqrt", IntrinsicType::kInverseSqrt},
+ IntrinsicData{"log", IntrinsicType::kLog},
+ IntrinsicData{"log2", IntrinsicType::kLog2},
+ IntrinsicData{"normalize", IntrinsicType::kNormalize},
+ IntrinsicData{"round", IntrinsicType::kRound},
+ IntrinsicData{"sign", IntrinsicType::kSign},
+ IntrinsicData{"sin", IntrinsicType::kSin},
+ IntrinsicData{"sinh", IntrinsicType::kSinh},
+ IntrinsicData{"sqrt", IntrinsicType::kSqrt},
+ IntrinsicData{"tan", IntrinsicType::kTan},
+ IntrinsicData{"tanh", IntrinsicType::kTanh},
+ IntrinsicData{"trunc", IntrinsicType::kTrunc}));
using ImportData_SingleParam_FloatOrInt_Test =
TypeDeterminerTestWithParam<IntrinsicData>;
@@ -1919,8 +1913,8 @@
INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest,
ImportData_SingleParam_FloatOrInt_Test,
- testing::Values(IntrinsicData{
- "abs", semantic::Intrinsic::kAbs}));
+ testing::Values(IntrinsicData{"abs",
+ IntrinsicType::kAbs}));
TEST_F(TypeDeterminerTest, ImportData_Length_Scalar) {
auto* ident = Expr("length");
@@ -1992,10 +1986,10 @@
INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest,
ImportData_TwoParamTest,
- testing::Values(IntrinsicData{"atan2", semantic::Intrinsic::kAtan2},
- IntrinsicData{"pow", semantic::Intrinsic::kPow},
- IntrinsicData{"step", semantic::Intrinsic::kStep},
- IntrinsicData{"reflect", semantic::Intrinsic::kReflect}));
+ testing::Values(IntrinsicData{"atan2", IntrinsicType::kAtan2},
+ IntrinsicData{"pow", IntrinsicType::kPow},
+ IntrinsicData{"step", IntrinsicType::kStep},
+ IntrinsicData{"reflect", IntrinsicType::kReflect}));
TEST_F(TypeDeterminerTest, ImportData_Distance_Scalar) {
auto* ident = Expr("distance");
@@ -2093,11 +2087,10 @@
INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest,
ImportData_ThreeParamTest,
- testing::Values(
- IntrinsicData{"mix", semantic::Intrinsic::kMix},
- IntrinsicData{"smoothStep", semantic::Intrinsic::kSmoothStep},
- IntrinsicData{"fma", semantic::Intrinsic::kFma},
- IntrinsicData{"faceForward", semantic::Intrinsic::kFaceForward}));
+ testing::Values(IntrinsicData{"mix", IntrinsicType::kMix},
+ IntrinsicData{"smoothStep", IntrinsicType::kSmoothStep},
+ IntrinsicData{"fma", IntrinsicType::kFma},
+ IntrinsicData{"faceForward", IntrinsicType::kFaceForward}));
using ImportData_ThreeParam_FloatOrInt_Test =
TypeDeterminerTestWithParam<IntrinsicData>;
@@ -2200,8 +2193,8 @@
INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest,
ImportData_ThreeParam_FloatOrInt_Test,
- testing::Values(IntrinsicData{
- "clamp", semantic::Intrinsic::kClamp}));
+ testing::Values(IntrinsicData{"clamp",
+ IntrinsicType::kClamp}));
using ImportData_Int_SingleParamTest =
TypeDeterminerTestWithParam<IntrinsicData>;
@@ -2248,9 +2241,8 @@
INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest,
ImportData_Int_SingleParamTest,
- testing::Values(
- IntrinsicData{"countOneBits", semantic::Intrinsic::kCountOneBits},
- IntrinsicData{"reverseBits", semantic::Intrinsic::kReverseBits}));
+ testing::Values(IntrinsicData{"countOneBits", IntrinsicType::kCountOneBits},
+ IntrinsicData{"reverseBits", IntrinsicType::kReverseBits}));
using ImportData_FloatOrInt_TwoParamTest =
TypeDeterminerTestWithParam<IntrinsicData>;
@@ -2351,8 +2343,8 @@
INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest,
ImportData_FloatOrInt_TwoParamTest,
- testing::Values(IntrinsicData{"min", semantic::Intrinsic::kMin},
- IntrinsicData{"max", semantic::Intrinsic::kMax}));
+ testing::Values(IntrinsicData{"min", IntrinsicType::kMin},
+ IntrinsicData{"max", IntrinsicType::kMax}));
TEST_F(TypeDeterminerTest, ImportData_GLSL_Determinant) {
Global("var", ast::StorageClass::kFunction, ty.mat3x3<f32>());
@@ -2385,8 +2377,7 @@
INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest,
ImportData_Matrix_OneParam_Test,
testing::Values(IntrinsicData{
- "determinant",
- semantic::Intrinsic::kDeterminant}));
+ "determinant", IntrinsicType::kDeterminant}));
TEST_F(TypeDeterminerTest, Function_EntryPoints_StageDecoration) {
// fn b() {}
diff --git a/src/validator/validator_impl.cc b/src/validator/validator_impl.cc
index bb568de..16da3e3 100644
--- a/src/validator/validator_impl.cc
+++ b/src/validator/validator_impl.cc
@@ -49,6 +49,8 @@
namespace {
+using IntrinsicType = semantic::IntrinsicType;
+
enum class IntrinsicDataType {
kMixed,
kFloatOrIntScalarOrVector,
@@ -63,7 +65,7 @@
};
struct IntrinsicData {
- semantic::Intrinsic intrinsic;
+ IntrinsicType intrinsic;
uint32_t param_count;
IntrinsicDataType data_type;
uint32_t vector_size;
@@ -73,122 +75,97 @@
// 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[] = {
- {semantic::Intrinsic::kAbs, 1, IntrinsicDataType::kFloatOrIntScalarOrVector,
- 0, true},
- {semantic::Intrinsic::kAcos, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kAbs, 1, IntrinsicDataType::kFloatOrIntScalarOrVector, 0,
true},
- {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,
+ {IntrinsicType::kAcos, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kAll, 1, IntrinsicDataType::kBoolVector, 0, false},
+ {IntrinsicType::kAny, 1, IntrinsicDataType::kBoolVector, 0, false},
+ {IntrinsicType::kArrayLength, 1, IntrinsicDataType::kMixed, 0, false},
+ {IntrinsicType::kAsin, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kAtan, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kAtan2, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
true},
- {semantic::Intrinsic::kAtan, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kCeil, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kClamp, 3, IntrinsicDataType::kFloatOrIntScalarOrVector, 0,
true},
- {semantic::Intrinsic::kAtan2, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kCos, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kCosh, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kCountOneBits, 1, IntrinsicDataType::kIntScalarOrVector, 0,
true},
- {semantic::Intrinsic::kCeil, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kClamp, 3,
- IntrinsicDataType::kFloatOrIntScalarOrVector, 0, true},
- {semantic::Intrinsic::kCos, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kCosh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kCountOneBits, 1,
- IntrinsicDataType::kIntScalarOrVector, 0, true},
- {semantic::Intrinsic::kCross, 2, IntrinsicDataType::kFloatVector, 3, true},
- {semantic::Intrinsic::kDeterminant, 1, IntrinsicDataType::kMatrix, 0,
+ {IntrinsicType::kCross, 2, IntrinsicDataType::kFloatVector, 3, true},
+ {IntrinsicType::kDeterminant, 1, IntrinsicDataType::kMatrix, 0, false},
+ {IntrinsicType::kDistance, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
false},
- {semantic::Intrinsic::kDistance, 2, IntrinsicDataType::kFloatScalarOrVector,
- 0, false},
- {semantic::Intrinsic::kDot, 2, IntrinsicDataType::kFloatVector, 0, false},
- {semantic::Intrinsic::kDpdx, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kDot, 2, IntrinsicDataType::kFloatVector, 0, false},
+ {IntrinsicType::kDpdx, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kDpdxCoarse, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true},
- {semantic::Intrinsic::kDpdxCoarse, 1,
- IntrinsicDataType::kFloatScalarOrVector, 0, true},
- {semantic::Intrinsic::kDpdxFine, 1, IntrinsicDataType::kFloatScalarOrVector,
+ {IntrinsicType::kDpdxFine, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ true},
+ {IntrinsicType::kDpdy, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kDpdyCoarse, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ true},
+ {IntrinsicType::kDpdyFine, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ true},
+ {IntrinsicType::kExp, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kExp2, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kFaceForward, 3, IntrinsicDataType::kFloatScalarOrVector, 0,
+ true},
+ {IntrinsicType::kFloor, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ true},
+ {IntrinsicType::kFma, 3, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kFract, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ true},
+ {IntrinsicType::kFrexp, 2, IntrinsicDataType::kMixed, 0, false},
+ {IntrinsicType::kFwidth, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ true},
+ {IntrinsicType::kFwidthCoarse, 1, IntrinsicDataType::kFloatScalarOrVector,
0, true},
- {semantic::Intrinsic::kDpdy, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kFwidthFine, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true},
- {semantic::Intrinsic::kDpdyCoarse, 1,
- IntrinsicDataType::kFloatScalarOrVector, 0, true},
- {semantic::Intrinsic::kDpdyFine, 1, IntrinsicDataType::kFloatScalarOrVector,
- 0, true},
- {semantic::Intrinsic::kExp, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kInverseSqrt, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true},
- {semantic::Intrinsic::kExp2, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kLdexp, 2, 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},
- {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},
- {semantic::Intrinsic::kLength, 1, IntrinsicDataType::kFloatScalarOrVector,
- 0, false},
- {semantic::Intrinsic::kLog, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kLog2, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kMax, 2, IntrinsicDataType::kFloatOrIntScalarOrVector,
- 0, true},
- {semantic::Intrinsic::kMin, 2, IntrinsicDataType::kFloatOrIntScalarOrVector,
- 0, true},
- {semantic::Intrinsic::kMix, 3, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kModf, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kNormalize, 1, IntrinsicDataType::kFloatVector, 0,
- true},
- {semantic::Intrinsic::kPack4x8Snorm, 1, IntrinsicDataType::kFloatVector, 4,
+ {IntrinsicType::kLength, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
false},
- {semantic::Intrinsic::kPack4x8Unorm, 1, IntrinsicDataType::kFloatVector, 4,
+ {IntrinsicType::kLog, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kLog2, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kMax, 2, IntrinsicDataType::kFloatOrIntScalarOrVector, 0,
+ true},
+ {IntrinsicType::kMin, 2, IntrinsicDataType::kFloatOrIntScalarOrVector, 0,
+ true},
+ {IntrinsicType::kMix, 3, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kModf, 2, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kNormalize, 1, IntrinsicDataType::kFloatVector, 0, true},
+ {IntrinsicType::kPack4x8Snorm, 1, IntrinsicDataType::kFloatVector, 4,
false},
- {semantic::Intrinsic::kPack2x16Snorm, 1, IntrinsicDataType::kFloatVector, 2,
+ {IntrinsicType::kPack4x8Unorm, 1, IntrinsicDataType::kFloatVector, 4,
false},
- {semantic::Intrinsic::kPack2x16Unorm, 1, IntrinsicDataType::kFloatVector, 2,
+ {IntrinsicType::kPack2x16Snorm, 1, IntrinsicDataType::kFloatVector, 2,
false},
- {semantic::Intrinsic::kPack2x16Float, 1, IntrinsicDataType::kFloatVector, 2,
+ {IntrinsicType::kPack2x16Unorm, 1, IntrinsicDataType::kFloatVector, 2,
false},
- {semantic::Intrinsic::kPow, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kPack2x16Float, 1, IntrinsicDataType::kFloatVector, 2,
+ false},
+ {IntrinsicType::kPow, 2, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kReflect, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
true},
- {semantic::Intrinsic::kReflect, 2, IntrinsicDataType::kFloatScalarOrVector,
- 0, true},
- {semantic::Intrinsic::kReverseBits, 1,
- IntrinsicDataType::kIntScalarOrVector, 0, true},
- {semantic::Intrinsic::kRound, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kReverseBits, 1, IntrinsicDataType::kIntScalarOrVector, 0,
true},
- {semantic::Intrinsic::kSelect, 3, IntrinsicDataType::kMixed, 0, false},
- {semantic::Intrinsic::kSign, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kRound, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true},
- {semantic::Intrinsic::kSin, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kSelect, 3, IntrinsicDataType::kMixed, 0, false},
+ {IntrinsicType::kSign, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kSin, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kSinh, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kSmoothStep, 3, IntrinsicDataType::kFloatScalarOrVector, 0,
true},
- {semantic::Intrinsic::kSinh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kSmoothStep, 3,
- IntrinsicDataType::kFloatScalarOrVector, 0, true},
- {semantic::Intrinsic::kSqrt, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kStep, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kTan, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kTanh, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
- true},
- {semantic::Intrinsic::kTrunc, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
+ {IntrinsicType::kSqrt, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kStep, 2, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kTan, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kTanh, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
+ {IntrinsicType::kTrunc, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true},
};
@@ -731,7 +708,7 @@
}
} else {
// Special cases.
- if (data->intrinsic == semantic::Intrinsic::kFrexp) {
+ if (data->intrinsic == IntrinsicType::kFrexp) {
auto* p0 = expr->params()[0];
auto* p1 = expr->params()[1];
auto* t0 = program_->TypeOf(p0)->UnwrapPtrIfNeeded();
@@ -773,7 +750,7 @@
}
}
- if (data->intrinsic == semantic::Intrinsic::kSelect) {
+ if (data->intrinsic == IntrinsicType::kSelect) {
auto* type = program_->TypeOf(expr);
auto* t0 = program_->TypeOf(expr->params()[0])->UnwrapPtrIfNeeded();
auto* t1 = program_->TypeOf(expr->params()[1])->UnwrapPtrIfNeeded();
@@ -819,7 +796,7 @@
}
}
- if (data->intrinsic == semantic::Intrinsic::kArrayLength) {
+ if (data->intrinsic == IntrinsicType::kArrayLength) {
if (!program_->TypeOf(expr)->UnwrapPtrIfNeeded()->Is<type::U32>()) {
add_error(
expr->source(),
@@ -840,15 +817,15 @@
}
// Result types don't match parameter types.
- if (data->intrinsic == semantic::Intrinsic::kAll ||
- data->intrinsic == semantic::Intrinsic::kAny) {
+ if (data->intrinsic == IntrinsicType::kAll ||
+ data->intrinsic == IntrinsicType::kAny) {
if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
IntrinsicDataType::kBoolScalar, 0, this)) {
return false;
}
}
- if (data->intrinsic == semantic::Intrinsic::kDot) {
+ if (data->intrinsic == IntrinsicType::kDot) {
if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
IntrinsicDataType::kFloatScalar, 0, this)) {
return false;
@@ -864,9 +841,9 @@
}
}
- if (data->intrinsic == semantic::Intrinsic::kLength ||
- data->intrinsic == semantic::Intrinsic::kDistance ||
- data->intrinsic == semantic::Intrinsic::kDeterminant) {
+ if (data->intrinsic == IntrinsicType::kLength ||
+ data->intrinsic == IntrinsicType::kDistance ||
+ data->intrinsic == IntrinsicType::kDeterminant) {
if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
IntrinsicDataType::kFloatScalar, 0, this)) {
return false;
@@ -874,7 +851,7 @@
}
// Must be a square matrix.
- if (data->intrinsic == semantic::Intrinsic::kDeterminant) {
+ if (data->intrinsic == IntrinsicType::kDeterminant) {
const auto* matrix =
program_->TypeOf(expr->params()[0])->As<type::Matrix>();
if (matrix->rows() != matrix->columns()) {
@@ -887,8 +864,8 @@
}
// Last parameter must be a pointer.
- if (data->intrinsic == semantic::Intrinsic::kFrexp ||
- data->intrinsic == semantic::Intrinsic::kModf) {
+ if (data->intrinsic == IntrinsicType::kFrexp ||
+ data->intrinsic == IntrinsicType::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 +
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 52dfa12..2a4d304 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -548,10 +548,10 @@
}
if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) {
const auto& params = expr->params();
- if (sem->intrinsic() == semantic::Intrinsic::kSelect) {
+ if (sem->intrinsic() == semantic::IntrinsicType::kSelect) {
error_ = "select not supported in HLSL backend yet";
return false;
- } else if (sem->intrinsic() == semantic::Intrinsic::kIsNormal) {
+ } else if (sem->intrinsic() == semantic::IntrinsicType::kIsNormal) {
error_ = "is_normal not supported in HLSL backend yet";
return false;
} else if (semantic::intrinsic::IsDataPackingIntrinsic(sem->intrinsic())) {
@@ -650,21 +650,21 @@
uint32_t dims = 2;
bool is_signed = false;
uint32_t scale = 65535;
- if (ident->intrinsic() == semantic::Intrinsic::kPack4x8Snorm ||
- ident->intrinsic() == semantic::Intrinsic::kPack4x8Unorm) {
+ if (ident->intrinsic() == semantic::IntrinsicType::kPack4x8Snorm ||
+ ident->intrinsic() == semantic::IntrinsicType::kPack4x8Unorm) {
dims = 4;
scale = 255;
}
- if (ident->intrinsic() == semantic::Intrinsic::kPack4x8Snorm ||
- ident->intrinsic() == semantic::Intrinsic::kPack2x16Snorm) {
+ if (ident->intrinsic() == semantic::IntrinsicType::kPack4x8Snorm ||
+ ident->intrinsic() == semantic::IntrinsicType::kPack2x16Snorm) {
is_signed = true;
scale = (scale - 1) / 2;
}
switch (ident->intrinsic()) {
- case semantic::Intrinsic::kPack4x8Snorm:
- case semantic::Intrinsic::kPack4x8Unorm:
- case semantic::Intrinsic::kPack2x16Snorm:
- case semantic::Intrinsic::kPack2x16Unorm:
+ case semantic::IntrinsicType::kPack4x8Snorm:
+ case semantic::IntrinsicType::kPack4x8Unorm:
+ case semantic::IntrinsicType::kPack2x16Snorm:
+ case semantic::IntrinsicType::kPack2x16Unorm:
pre << (is_signed ? "" : "u") << "int" << dims << " " << tmp_name << " = "
<< (is_signed ? "" : "u") << "int" << dims << "(round(clamp("
<< expr_out.str() << ", " << (is_signed ? "-1.0" : "0.0")
@@ -683,7 +683,7 @@
}
out << ")";
break;
- case semantic::Intrinsic::kPack2x16Float:
+ case semantic::IntrinsicType::kPack2x16Float:
pre << "uint2 " << tmp_name << " = f32tof16(" << expr_out.str() << ");\n";
out << "(" << tmp_name << ".x | " << tmp_name << ".y << 16)";
break;
@@ -709,17 +709,17 @@
auto* texture_type = TypeOf(texture)->UnwrapAll()->As<type::Texture>();
switch (sem->intrinsic()) {
- case semantic::Intrinsic::kTextureDimensions:
- case semantic::Intrinsic::kTextureNumLayers:
- case semantic::Intrinsic::kTextureNumLevels:
- case semantic::Intrinsic::kTextureNumSamples: {
+ case semantic::IntrinsicType::kTextureDimensions:
+ case semantic::IntrinsicType::kTextureNumLayers:
+ case semantic::IntrinsicType::kTextureNumLevels:
+ case semantic::IntrinsicType::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 (sem->intrinsic()) {
- case semantic::Intrinsic::kTextureDimensions:
+ case semantic::IntrinsicType::kTextureDimensions:
switch (texture_type->dim()) {
case type::TextureDimension::kNone:
error_ = "texture dimension is kNone";
@@ -755,7 +755,7 @@
break;
}
break;
- case semantic::Intrinsic::kTextureNumLayers:
+ case semantic::IntrinsicType::kTextureNumLayers:
switch (texture_type->dim()) {
default:
error_ = "texture dimension is not arrayed";
@@ -771,7 +771,7 @@
break;
}
break;
- case semantic::Intrinsic::kTextureNumLevels:
+ case semantic::IntrinsicType::kTextureNumLevels:
add_mip_level_in = true;
switch (texture_type->dim()) {
default:
@@ -790,7 +790,7 @@
break;
}
break;
- case semantic::Intrinsic::kTextureNumSamples:
+ case semantic::IntrinsicType::kTextureNumSamples:
switch (texture_type->dim()) {
default:
error_ = "texture dimension does not support multisampling";
@@ -860,28 +860,28 @@
bool pack_mip_in_coords = false;
switch (sem->intrinsic()) {
- case semantic::Intrinsic::kTextureSample:
+ case semantic::IntrinsicType::kTextureSample:
out << ".Sample(";
break;
- case semantic::Intrinsic::kTextureSampleBias:
+ case semantic::IntrinsicType::kTextureSampleBias:
out << ".SampleBias(";
break;
- case semantic::Intrinsic::kTextureSampleLevel:
+ case semantic::IntrinsicType::kTextureSampleLevel:
out << ".SampleLevel(";
break;
- case semantic::Intrinsic::kTextureSampleGrad:
+ case semantic::IntrinsicType::kTextureSampleGrad:
out << ".SampleGrad(";
break;
- case semantic::Intrinsic::kTextureSampleCompare:
+ case semantic::IntrinsicType::kTextureSampleCompare:
out << ".SampleCmp(";
break;
- case semantic::Intrinsic::kTextureLoad:
+ case semantic::IntrinsicType::kTextureLoad:
out << ".Load(";
if (!texture_type->Is<type::StorageTexture>()) {
pack_mip_in_coords = true;
}
break;
- case semantic::Intrinsic::kTextureStore:
+ case semantic::IntrinsicType::kTextureStore:
out << "[";
break;
default:
@@ -937,7 +937,7 @@
}
}
- if (sem->intrinsic() == semantic::Intrinsic::kTextureStore) {
+ if (sem->intrinsic() == semantic::IntrinsicType::kTextureStore) {
out << "] = ";
if (!EmitExpression(pre, out, params[pidx.value]))
return false;
@@ -952,94 +952,94 @@
const semantic::IntrinsicCall* call) {
std::string out;
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:
+ case semantic::IntrinsicType::kAcos:
+ case semantic::IntrinsicType::kAny:
+ case semantic::IntrinsicType::kAll:
+ case semantic::IntrinsicType::kAsin:
+ case semantic::IntrinsicType::kAtan:
+ case semantic::IntrinsicType::kAtan2:
+ case semantic::IntrinsicType::kCeil:
+ case semantic::IntrinsicType::kCos:
+ case semantic::IntrinsicType::kCosh:
+ case semantic::IntrinsicType::kCross:
+ case semantic::IntrinsicType::kDeterminant:
+ case semantic::IntrinsicType::kDistance:
+ case semantic::IntrinsicType::kDot:
+ case semantic::IntrinsicType::kExp:
+ case semantic::IntrinsicType::kExp2:
+ case semantic::IntrinsicType::kFloor:
+ case semantic::IntrinsicType::kFma:
+ case semantic::IntrinsicType::kLdexp:
+ case semantic::IntrinsicType::kLength:
+ case semantic::IntrinsicType::kLog:
+ case semantic::IntrinsicType::kLog2:
+ case semantic::IntrinsicType::kNormalize:
+ case semantic::IntrinsicType::kPow:
+ case semantic::IntrinsicType::kReflect:
+ case semantic::IntrinsicType::kRound:
+ case semantic::IntrinsicType::kSin:
+ case semantic::IntrinsicType::kSinh:
+ case semantic::IntrinsicType::kSqrt:
+ case semantic::IntrinsicType::kStep:
+ case semantic::IntrinsicType::kTan:
+ case semantic::IntrinsicType::kTanh:
+ case semantic::IntrinsicType::kTrunc:
+ case semantic::IntrinsicType::kMix:
+ case semantic::IntrinsicType::kSign:
+ case semantic::IntrinsicType::kAbs:
+ case semantic::IntrinsicType::kMax:
+ case semantic::IntrinsicType::kMin:
+ case semantic::IntrinsicType::kClamp:
out = semantic::intrinsic::str(call->intrinsic());
break;
- case semantic::Intrinsic::kCountOneBits:
+ case semantic::IntrinsicType::kCountOneBits:
out = "countbits";
break;
- case semantic::Intrinsic::kDpdx:
+ case semantic::IntrinsicType::kDpdx:
out = "ddx";
break;
- case semantic::Intrinsic::kDpdxCoarse:
+ case semantic::IntrinsicType::kDpdxCoarse:
out = "ddx_coarse";
break;
- case semantic::Intrinsic::kDpdxFine:
+ case semantic::IntrinsicType::kDpdxFine:
out = "ddx_fine";
break;
- case semantic::Intrinsic::kDpdy:
+ case semantic::IntrinsicType::kDpdy:
out = "ddy";
break;
- case semantic::Intrinsic::kDpdyCoarse:
+ case semantic::IntrinsicType::kDpdyCoarse:
out = "ddy_coarse";
break;
- case semantic::Intrinsic::kDpdyFine:
+ case semantic::IntrinsicType::kDpdyFine:
out = "ddy_fine";
break;
- case semantic::Intrinsic::kFaceForward:
+ case semantic::IntrinsicType::kFaceForward:
out = "faceforward";
break;
- case semantic::Intrinsic::kFract:
+ case semantic::IntrinsicType::kFract:
out = "frac";
break;
- case semantic::Intrinsic::kFwidth:
- case semantic::Intrinsic::kFwidthCoarse:
- case semantic::Intrinsic::kFwidthFine:
+ case semantic::IntrinsicType::kFwidth:
+ case semantic::IntrinsicType::kFwidthCoarse:
+ case semantic::IntrinsicType::kFwidthFine:
out = "fwidth";
break;
- case semantic::Intrinsic::kInverseSqrt:
+ case semantic::IntrinsicType::kInverseSqrt:
out = "rsqrt";
break;
- case semantic::Intrinsic::kIsFinite:
+ case semantic::IntrinsicType::kIsFinite:
out = "isfinite";
break;
- case semantic::Intrinsic::kIsInf:
+ case semantic::IntrinsicType::kIsInf:
out = "isinf";
break;
- case semantic::Intrinsic::kIsNan:
+ case semantic::IntrinsicType::kIsNan:
out = "isnan";
break;
- case semantic::Intrinsic::kReverseBits:
+ case semantic::IntrinsicType::kReverseBits:
out = "reversebits";
break;
- case semantic::Intrinsic::kSmoothStep:
+ case semantic::IntrinsicType::kSmoothStep:
out = "smoothstep";
break;
default:
diff --git a/src/writer/hlsl/generator_impl_intrinsic_test.cc b/src/writer/hlsl/generator_impl_intrinsic_test.cc
index 1976462..27a76e5 100644
--- a/src/writer/hlsl/generator_impl_intrinsic_test.cc
+++ b/src/writer/hlsl/generator_impl_intrinsic_test.cc
@@ -29,6 +29,8 @@
namespace hlsl {
namespace {
+using IntrinsicType = semantic::IntrinsicType;
+
using ::testing::HasSubstr;
using HlslGeneratorImplTest_Intrinsic = TestHelper;
@@ -40,7 +42,7 @@
};
struct IntrinsicData {
- semantic::Intrinsic intrinsic;
+ IntrinsicType intrinsic;
ParamType type;
const char* hlsl_name;
};
@@ -61,92 +63,92 @@
return out;
}
-ast::CallExpression* GenerateCall(semantic::Intrinsic intrinsic,
+ast::CallExpression* GenerateCall(IntrinsicType intrinsic,
ParamType type,
ProgramBuilder* builder) {
std::string name;
std::ostringstream str(name);
str << intrinsic;
switch (intrinsic) {
- 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:
+ case IntrinsicType::kAcos:
+ case IntrinsicType::kAsin:
+ case IntrinsicType::kAtan:
+ case IntrinsicType::kCeil:
+ case IntrinsicType::kCos:
+ case IntrinsicType::kCosh:
+ case IntrinsicType::kDpdx:
+ case IntrinsicType::kDpdxCoarse:
+ case IntrinsicType::kDpdxFine:
+ case IntrinsicType::kDpdy:
+ case IntrinsicType::kDpdyCoarse:
+ case IntrinsicType::kDpdyFine:
+ case IntrinsicType::kExp:
+ case IntrinsicType::kExp2:
+ case IntrinsicType::kFloor:
+ case IntrinsicType::kFract:
+ case IntrinsicType::kFwidth:
+ case IntrinsicType::kFwidthCoarse:
+ case IntrinsicType::kFwidthFine:
+ case IntrinsicType::kInverseSqrt:
+ case IntrinsicType::kIsFinite:
+ case IntrinsicType::kIsInf:
+ case IntrinsicType::kIsNan:
+ case IntrinsicType::kIsNormal:
+ case IntrinsicType::kLdexp:
+ case IntrinsicType::kLength:
+ case IntrinsicType::kLog:
+ case IntrinsicType::kLog2:
+ case IntrinsicType::kNormalize:
+ case IntrinsicType::kReflect:
+ case IntrinsicType::kRound:
+ case IntrinsicType::kSin:
+ case IntrinsicType::kSinh:
+ case IntrinsicType::kSqrt:
+ case IntrinsicType::kTan:
+ case IntrinsicType::kTanh:
+ case IntrinsicType::kTrunc:
+ case IntrinsicType::kSign:
return builder->Call(str.str(), "f1");
- case semantic::Intrinsic::kAtan2:
- case semantic::Intrinsic::kCross:
- case semantic::Intrinsic::kDot:
- case semantic::Intrinsic::kDistance:
- case semantic::Intrinsic::kPow:
- case semantic::Intrinsic::kStep:
+ case IntrinsicType::kAtan2:
+ case IntrinsicType::kCross:
+ case IntrinsicType::kDot:
+ case IntrinsicType::kDistance:
+ case IntrinsicType::kPow:
+ case IntrinsicType::kStep:
return builder->Call(str.str(), "f1", "f2");
- case semantic::Intrinsic::kFma:
- case semantic::Intrinsic::kMix:
- case semantic::Intrinsic::kFaceForward:
- case semantic::Intrinsic::kSmoothStep:
+ case IntrinsicType::kFma:
+ case IntrinsicType::kMix:
+ case IntrinsicType::kFaceForward:
+ case IntrinsicType::kSmoothStep:
return builder->Call(str.str(), "f1", "f2", "f3");
- case semantic::Intrinsic::kAll:
- case semantic::Intrinsic::kAny:
+ case IntrinsicType::kAll:
+ case IntrinsicType::kAny:
return builder->Call(str.str(), "b1");
- case semantic::Intrinsic::kAbs:
+ case IntrinsicType::kAbs:
if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1");
} else {
return builder->Call(str.str(), "u1");
}
- case semantic::Intrinsic::kCountOneBits:
- case semantic::Intrinsic::kReverseBits:
+ case IntrinsicType::kCountOneBits:
+ case IntrinsicType::kReverseBits:
return builder->Call(str.str(), "u1");
- case semantic::Intrinsic::kMax:
- case semantic::Intrinsic::kMin:
+ case IntrinsicType::kMax:
+ case IntrinsicType::kMin:
if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1", "f2");
} else {
return builder->Call(str.str(), "u1", "u2");
}
- case semantic::Intrinsic::kClamp:
+ case IntrinsicType::kClamp:
if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1", "f2", "f3");
} else {
return builder->Call(str.str(), "u1", "u2", "u3");
}
- case semantic::Intrinsic::kSelect:
+ case IntrinsicType::kSelect:
return builder->Call(str.str(), "f1", "f2", "b1");
- case semantic::Intrinsic::kDeterminant:
+ case IntrinsicType::kDeterminant:
return builder->Call(str.str(), "m1");
default:
break;
@@ -183,81 +185,72 @@
HlslGeneratorImplTest_Intrinsic,
HlslIntrinsicTest,
testing::Values(
- 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,
+ IntrinsicData{IntrinsicType::kAbs, ParamType::kF32, "abs"},
+ IntrinsicData{IntrinsicType::kAbs, ParamType::kU32, "abs"},
+ IntrinsicData{IntrinsicType::kAcos, ParamType::kF32, "acos"},
+ IntrinsicData{IntrinsicType::kAll, ParamType::kBool, "all"},
+ IntrinsicData{IntrinsicType::kAny, ParamType::kBool, "any"},
+ IntrinsicData{IntrinsicType::kAsin, ParamType::kF32, "asin"},
+ IntrinsicData{IntrinsicType::kAtan, ParamType::kF32, "atan"},
+ IntrinsicData{IntrinsicType::kAtan2, ParamType::kF32, "atan2"},
+ IntrinsicData{IntrinsicType::kCeil, ParamType::kF32, "ceil"},
+ IntrinsicData{IntrinsicType::kClamp, ParamType::kF32, "clamp"},
+ IntrinsicData{IntrinsicType::kClamp, ParamType::kU32, "clamp"},
+ IntrinsicData{IntrinsicType::kCos, ParamType::kF32, "cos"},
+ IntrinsicData{IntrinsicType::kCosh, ParamType::kF32, "cosh"},
+ IntrinsicData{IntrinsicType::kCountOneBits, ParamType::kU32,
"countbits"},
- IntrinsicData{semantic::Intrinsic::kCross, ParamType::kF32, "cross"},
- IntrinsicData{semantic::Intrinsic::kDeterminant, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kCross, ParamType::kF32, "cross"},
+ IntrinsicData{IntrinsicType::kDeterminant, ParamType::kF32,
"determinant"},
- 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,
+ IntrinsicData{IntrinsicType::kDistance, ParamType::kF32, "distance"},
+ IntrinsicData{IntrinsicType::kDot, ParamType::kF32, "dot"},
+ IntrinsicData{IntrinsicType::kDpdx, ParamType::kF32, "ddx"},
+ IntrinsicData{IntrinsicType::kDpdxCoarse, ParamType::kF32,
"ddx_coarse"},
- IntrinsicData{semantic::Intrinsic::kDpdxFine, ParamType::kF32,
- "ddx_fine"},
- IntrinsicData{semantic::Intrinsic::kDpdy, ParamType::kF32, "ddy"},
- IntrinsicData{semantic::Intrinsic::kDpdyCoarse, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kDpdxFine, ParamType::kF32, "ddx_fine"},
+ IntrinsicData{IntrinsicType::kDpdy, ParamType::kF32, "ddy"},
+ IntrinsicData{IntrinsicType::kDpdyCoarse, ParamType::kF32,
"ddy_coarse"},
- 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,
+ IntrinsicData{IntrinsicType::kDpdyFine, ParamType::kF32, "ddy_fine"},
+ IntrinsicData{IntrinsicType::kExp, ParamType::kF32, "exp"},
+ IntrinsicData{IntrinsicType::kExp2, ParamType::kF32, "exp2"},
+ IntrinsicData{IntrinsicType::kFaceForward, ParamType::kF32,
"faceforward"},
- 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,
+ IntrinsicData{IntrinsicType::kFloor, ParamType::kF32, "floor"},
+ IntrinsicData{IntrinsicType::kFma, ParamType::kF32, "fma"},
+ IntrinsicData{IntrinsicType::kFract, ParamType::kF32, "frac"},
+ IntrinsicData{IntrinsicType::kFwidth, ParamType::kF32, "fwidth"},
+ IntrinsicData{IntrinsicType::kFwidthCoarse, ParamType::kF32, "fwidth"},
+ IntrinsicData{IntrinsicType::kFwidthFine, ParamType::kF32, "fwidth"},
+ IntrinsicData{IntrinsicType::kInverseSqrt, ParamType::kF32, "rsqrt"},
+ IntrinsicData{IntrinsicType::kIsFinite, ParamType::kF32, "isfinite"},
+ IntrinsicData{IntrinsicType::kIsInf, ParamType::kF32, "isinf"},
+ IntrinsicData{IntrinsicType::kIsNan, ParamType::kF32, "isnan"},
+ IntrinsicData{IntrinsicType::kLdexp, ParamType::kF32, "ldexp"},
+ IntrinsicData{IntrinsicType::kLength, ParamType::kF32, "length"},
+ IntrinsicData{IntrinsicType::kLog, ParamType::kF32, "log"},
+ IntrinsicData{IntrinsicType::kLog2, ParamType::kF32, "log2"},
+ IntrinsicData{IntrinsicType::kMax, ParamType::kF32, "max"},
+ IntrinsicData{IntrinsicType::kMax, ParamType::kU32, "max"},
+ IntrinsicData{IntrinsicType::kMin, ParamType::kF32, "min"},
+ IntrinsicData{IntrinsicType::kMin, ParamType::kU32, "min"},
+ IntrinsicData{IntrinsicType::kNormalize, ParamType::kF32, "normalize"},
+ IntrinsicData{IntrinsicType::kPow, ParamType::kF32, "pow"},
+ IntrinsicData{IntrinsicType::kReflect, ParamType::kF32, "reflect"},
+ IntrinsicData{IntrinsicType::kReverseBits, ParamType::kU32,
"reversebits"},
- 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,
+ IntrinsicData{IntrinsicType::kRound, ParamType::kU32, "round"},
+ IntrinsicData{IntrinsicType::kSign, ParamType::kF32, "sign"},
+ IntrinsicData{IntrinsicType::kSin, ParamType::kF32, "sin"},
+ IntrinsicData{IntrinsicType::kSinh, ParamType::kF32, "sinh"},
+ IntrinsicData{IntrinsicType::kSmoothStep, ParamType::kF32,
"smoothstep"},
- 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"}));
+ IntrinsicData{IntrinsicType::kSqrt, ParamType::kF32, "sqrt"},
+ IntrinsicData{IntrinsicType::kStep, ParamType::kF32, "step"},
+ IntrinsicData{IntrinsicType::kTan, ParamType::kF32, "tan"},
+ IntrinsicData{IntrinsicType::kTanh, ParamType::kF32, "tanh"},
+ IntrinsicData{IntrinsicType::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 cdce707..0020501 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -449,7 +449,7 @@
return EmitTextureCall(expr, sem);
}
if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) {
- if (sem->intrinsic() == semantic::Intrinsic::kPack2x16Float) {
+ if (sem->intrinsic() == semantic::IntrinsicType::kPack2x16Float) {
make_indent();
out_ << "as_type<uint>(half2(";
if (!EmitExpression(expr->params()[0])) {
@@ -580,7 +580,7 @@
TypeOf(params[pidx.texture])->UnwrapAll()->As<type::Texture>();
switch (sem->intrinsic()) {
- case semantic::Intrinsic::kTextureDimensions: {
+ case semantic::IntrinsicType::kTextureDimensions: {
std::vector<const char*> dims;
switch (texture_type->dim()) {
case type::TextureDimension::kNone:
@@ -634,7 +634,7 @@
}
return true;
}
- case semantic::Intrinsic::kTextureNumLayers: {
+ case semantic::IntrinsicType::kTextureNumLayers: {
out_ << "int(";
if (!EmitExpression(params[pidx.texture])) {
return false;
@@ -642,7 +642,7 @@
out_ << ".get_array_size())";
return true;
}
- case semantic::Intrinsic::kTextureNumLevels: {
+ case semantic::IntrinsicType::kTextureNumLevels: {
out_ << "int(";
if (!EmitExpression(params[pidx.texture])) {
return false;
@@ -650,7 +650,7 @@
out_ << ".get_num_mip_levels())";
return true;
}
- case semantic::Intrinsic::kTextureNumSamples: {
+ case semantic::IntrinsicType::kTextureNumSamples: {
out_ << "int(";
if (!EmitExpression(params[pidx.texture])) {
return false;
@@ -668,20 +668,20 @@
bool lod_param_is_named = true;
switch (sem->intrinsic()) {
- case semantic::Intrinsic::kTextureSample:
- case semantic::Intrinsic::kTextureSampleBias:
- case semantic::Intrinsic::kTextureSampleLevel:
- case semantic::Intrinsic::kTextureSampleGrad:
+ case semantic::IntrinsicType::kTextureSample:
+ case semantic::IntrinsicType::kTextureSampleBias:
+ case semantic::IntrinsicType::kTextureSampleLevel:
+ case semantic::IntrinsicType::kTextureSampleGrad:
out_ << ".sample(";
break;
- case semantic::Intrinsic::kTextureSampleCompare:
+ case semantic::IntrinsicType::kTextureSampleCompare:
out_ << ".sample_compare(";
break;
- case semantic::Intrinsic::kTextureLoad:
+ case semantic::IntrinsicType::kTextureLoad:
out_ << ".read(";
lod_param_is_named = false;
break;
- case semantic::Intrinsic::kTextureStore:
+ case semantic::IntrinsicType::kTextureStore:
out_ << ".write(";
break;
default:
@@ -780,118 +780,118 @@
const semantic::IntrinsicCall* call) {
std::string out = "metal::";
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:
+ case semantic::IntrinsicType::kAcos:
+ case semantic::IntrinsicType::kAll:
+ case semantic::IntrinsicType::kAny:
+ case semantic::IntrinsicType::kAsin:
+ case semantic::IntrinsicType::kAtan:
+ case semantic::IntrinsicType::kAtan2:
+ case semantic::IntrinsicType::kCeil:
+ case semantic::IntrinsicType::kCos:
+ case semantic::IntrinsicType::kCosh:
+ case semantic::IntrinsicType::kCross:
+ case semantic::IntrinsicType::kDeterminant:
+ case semantic::IntrinsicType::kDistance:
+ case semantic::IntrinsicType::kDot:
+ case semantic::IntrinsicType::kExp:
+ case semantic::IntrinsicType::kExp2:
+ case semantic::IntrinsicType::kFloor:
+ case semantic::IntrinsicType::kFma:
+ case semantic::IntrinsicType::kFract:
+ case semantic::IntrinsicType::kLength:
+ case semantic::IntrinsicType::kLdexp:
+ case semantic::IntrinsicType::kLog:
+ case semantic::IntrinsicType::kLog2:
+ case semantic::IntrinsicType::kMix:
+ case semantic::IntrinsicType::kNormalize:
+ case semantic::IntrinsicType::kPow:
+ case semantic::IntrinsicType::kReflect:
+ case semantic::IntrinsicType::kRound:
+ case semantic::IntrinsicType::kSelect:
+ case semantic::IntrinsicType::kSin:
+ case semantic::IntrinsicType::kSinh:
+ case semantic::IntrinsicType::kSqrt:
+ case semantic::IntrinsicType::kStep:
+ case semantic::IntrinsicType::kTan:
+ case semantic::IntrinsicType::kTanh:
+ case semantic::IntrinsicType::kTrunc:
+ case semantic::IntrinsicType::kSign:
+ case semantic::IntrinsicType::kClamp:
out += semantic::intrinsic::str(call->intrinsic());
break;
- case semantic::Intrinsic::kAbs:
+ case semantic::IntrinsicType::kAbs:
if (call->Type()->is_float_scalar_or_vector()) {
out += "fabs";
} else {
out += "abs";
}
break;
- case semantic::Intrinsic::kCountOneBits:
+ case semantic::IntrinsicType::kCountOneBits:
out += "popcount";
break;
- case semantic::Intrinsic::kDpdx:
- case semantic::Intrinsic::kDpdxCoarse:
- case semantic::Intrinsic::kDpdxFine:
+ case semantic::IntrinsicType::kDpdx:
+ case semantic::IntrinsicType::kDpdxCoarse:
+ case semantic::IntrinsicType::kDpdxFine:
out += "dfdx";
break;
- case semantic::Intrinsic::kDpdy:
- case semantic::Intrinsic::kDpdyCoarse:
- case semantic::Intrinsic::kDpdyFine:
+ case semantic::IntrinsicType::kDpdy:
+ case semantic::IntrinsicType::kDpdyCoarse:
+ case semantic::IntrinsicType::kDpdyFine:
out += "dfdy";
break;
- case semantic::Intrinsic::kFwidth:
- case semantic::Intrinsic::kFwidthCoarse:
- case semantic::Intrinsic::kFwidthFine:
+ case semantic::IntrinsicType::kFwidth:
+ case semantic::IntrinsicType::kFwidthCoarse:
+ case semantic::IntrinsicType::kFwidthFine:
out += "fwidth";
break;
- case semantic::Intrinsic::kIsFinite:
+ case semantic::IntrinsicType::kIsFinite:
out += "isfinite";
break;
- case semantic::Intrinsic::kIsInf:
+ case semantic::IntrinsicType::kIsInf:
out += "isinf";
break;
- case semantic::Intrinsic::kIsNan:
+ case semantic::IntrinsicType::kIsNan:
out += "isnan";
break;
- case semantic::Intrinsic::kIsNormal:
+ case semantic::IntrinsicType::kIsNormal:
out += "isnormal";
break;
- case semantic::Intrinsic::kMax:
+ case semantic::IntrinsicType::kMax:
if (call->Type()->is_float_scalar_or_vector()) {
out += "fmax";
} else {
out += "max";
}
break;
- case semantic::Intrinsic::kMin:
+ case semantic::IntrinsicType::kMin:
if (call->Type()->is_float_scalar_or_vector()) {
out += "fmin";
} else {
out += "min";
}
break;
- case semantic::Intrinsic::kFaceForward:
+ case semantic::IntrinsicType::kFaceForward:
out += "faceforward";
break;
- case semantic::Intrinsic::kPack4x8Snorm:
+ case semantic::IntrinsicType::kPack4x8Snorm:
out += "pack_float_to_snorm4x8";
break;
- case semantic::Intrinsic::kPack4x8Unorm:
+ case semantic::IntrinsicType::kPack4x8Unorm:
out += "pack_float_to_unorm4x8";
break;
- case semantic::Intrinsic::kPack2x16Snorm:
+ case semantic::IntrinsicType::kPack2x16Snorm:
out += "pack_float_to_snorm2x16";
break;
- case semantic::Intrinsic::kPack2x16Unorm:
+ case semantic::IntrinsicType::kPack2x16Unorm:
out += "pack_float_to_unorm2x16";
break;
- case semantic::Intrinsic::kReverseBits:
+ case semantic::IntrinsicType::kReverseBits:
out += "reverse_bits";
break;
- case semantic::Intrinsic::kSmoothStep:
+ case semantic::IntrinsicType::kSmoothStep:
out += "smoothstep";
break;
- case semantic::Intrinsic::kInverseSqrt:
+ case semantic::IntrinsicType::kInverseSqrt:
out += "rsqrt";
break;
default:
diff --git a/src/writer/msl/generator_impl_intrinsic_test.cc b/src/writer/msl/generator_impl_intrinsic_test.cc
index 7540043..2555904 100644
--- a/src/writer/msl/generator_impl_intrinsic_test.cc
+++ b/src/writer/msl/generator_impl_intrinsic_test.cc
@@ -30,6 +30,8 @@
namespace msl {
namespace {
+using IntrinsicType = semantic::IntrinsicType;
+
using MslGeneratorImplTest = TestHelper;
enum class ParamType {
@@ -39,7 +41,7 @@
};
struct IntrinsicData {
- semantic::Intrinsic intrinsic;
+ IntrinsicType intrinsic;
ParamType type;
const char* msl_name;
};
@@ -60,97 +62,97 @@
return out;
}
-ast::CallExpression* GenerateCall(semantic::Intrinsic intrinsic,
+ast::CallExpression* GenerateCall(IntrinsicType intrinsic,
ParamType type,
ProgramBuilder* builder) {
std::string name;
std::ostringstream str(name);
str << intrinsic;
switch (intrinsic) {
- 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::kPack4x8Snorm:
- case semantic::Intrinsic::kPack4x8Unorm:
- 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:
+ case IntrinsicType::kAcos:
+ case IntrinsicType::kAsin:
+ case IntrinsicType::kAtan:
+ case IntrinsicType::kCeil:
+ case IntrinsicType::kCos:
+ case IntrinsicType::kCosh:
+ case IntrinsicType::kDpdx:
+ case IntrinsicType::kDpdxCoarse:
+ case IntrinsicType::kDpdxFine:
+ case IntrinsicType::kDpdy:
+ case IntrinsicType::kDpdyCoarse:
+ case IntrinsicType::kDpdyFine:
+ case IntrinsicType::kExp:
+ case IntrinsicType::kExp2:
+ case IntrinsicType::kFloor:
+ case IntrinsicType::kFract:
+ case IntrinsicType::kFwidth:
+ case IntrinsicType::kFwidthCoarse:
+ case IntrinsicType::kFwidthFine:
+ case IntrinsicType::kInverseSqrt:
+ case IntrinsicType::kIsFinite:
+ case IntrinsicType::kIsInf:
+ case IntrinsicType::kIsNan:
+ case IntrinsicType::kIsNormal:
+ case IntrinsicType::kLdexp:
+ case IntrinsicType::kLength:
+ case IntrinsicType::kLog:
+ case IntrinsicType::kLog2:
+ case IntrinsicType::kNormalize:
+ case IntrinsicType::kPack4x8Snorm:
+ case IntrinsicType::kPack4x8Unorm:
+ case IntrinsicType::kReflect:
+ case IntrinsicType::kRound:
+ case IntrinsicType::kSin:
+ case IntrinsicType::kSinh:
+ case IntrinsicType::kSqrt:
+ case IntrinsicType::kTan:
+ case IntrinsicType::kTanh:
+ case IntrinsicType::kTrunc:
+ case IntrinsicType::kSign:
return builder->Call(str.str(), "f1");
- case semantic::Intrinsic::kAtan2:
- case semantic::Intrinsic::kCross:
- case semantic::Intrinsic::kDot:
- case semantic::Intrinsic::kDistance:
- case semantic::Intrinsic::kPow:
- case semantic::Intrinsic::kStep:
+ case IntrinsicType::kAtan2:
+ case IntrinsicType::kCross:
+ case IntrinsicType::kDot:
+ case IntrinsicType::kDistance:
+ case IntrinsicType::kPow:
+ case IntrinsicType::kStep:
return builder->Call(str.str(), "f1", "f2");
- case semantic::Intrinsic::kFma:
- case semantic::Intrinsic::kMix:
- case semantic::Intrinsic::kFaceForward:
- case semantic::Intrinsic::kSmoothStep:
+ case IntrinsicType::kFma:
+ case IntrinsicType::kMix:
+ case IntrinsicType::kFaceForward:
+ case IntrinsicType::kSmoothStep:
return builder->Call(str.str(), "f1", "f2", "f3");
- case semantic::Intrinsic::kAll:
- case semantic::Intrinsic::kAny:
+ case IntrinsicType::kAll:
+ case IntrinsicType::kAny:
return builder->Call(str.str(), "b1");
- case semantic::Intrinsic::kAbs:
+ case IntrinsicType::kAbs:
if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1");
} else {
return builder->Call(str.str(), "u1");
}
- case semantic::Intrinsic::kCountOneBits:
- case semantic::Intrinsic::kReverseBits:
+ case IntrinsicType::kCountOneBits:
+ case IntrinsicType::kReverseBits:
return builder->Call(str.str(), "u1");
- case semantic::Intrinsic::kMax:
- case semantic::Intrinsic::kMin:
+ case IntrinsicType::kMax:
+ case IntrinsicType::kMin:
if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1", "f2");
} else {
return builder->Call(str.str(), "u1", "u2");
}
- case semantic::Intrinsic::kClamp:
+ case IntrinsicType::kClamp:
if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1", "f2", "f3");
} else {
return builder->Call(str.str(), "u1", "u2", "u3");
}
- case semantic::Intrinsic::kSelect:
+ case IntrinsicType::kSelect:
return builder->Call(str.str(), "f1", "f2", "b1");
- case semantic::Intrinsic::kDeterminant:
+ case IntrinsicType::kDeterminant:
return builder->Call(str.str(), "m1");
- case semantic::Intrinsic::kPack2x16Snorm:
- case semantic::Intrinsic::kPack2x16Unorm:
+ case IntrinsicType::kPack2x16Snorm:
+ case IntrinsicType::kPack2x16Unorm:
return builder->Call(str.str(), "f4");
default:
break;
@@ -189,125 +191,90 @@
MslGeneratorImplTest,
MslIntrinsicTest,
testing::Values(
- 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,
+ IntrinsicData{IntrinsicType::kAbs, ParamType::kF32, "metal::fabs"},
+ IntrinsicData{IntrinsicType::kAbs, ParamType::kU32, "metal::abs"},
+ IntrinsicData{IntrinsicType::kAcos, ParamType::kF32, "metal::acos"},
+ IntrinsicData{IntrinsicType::kAll, ParamType::kBool, "metal::all"},
+ IntrinsicData{IntrinsicType::kAny, ParamType::kBool, "metal::any"},
+ IntrinsicData{IntrinsicType::kAsin, ParamType::kF32, "metal::asin"},
+ IntrinsicData{IntrinsicType::kAtan, ParamType::kF32, "metal::atan"},
+ IntrinsicData{IntrinsicType::kAtan2, ParamType::kF32, "metal::atan2"},
+ IntrinsicData{IntrinsicType::kCeil, ParamType::kF32, "metal::ceil"},
+ IntrinsicData{IntrinsicType::kClamp, ParamType::kF32, "metal::clamp"},
+ IntrinsicData{IntrinsicType::kClamp, ParamType::kU32, "metal::clamp"},
+ IntrinsicData{IntrinsicType::kCos, ParamType::kF32, "metal::cos"},
+ IntrinsicData{IntrinsicType::kCosh, ParamType::kF32, "metal::cosh"},
+ IntrinsicData{IntrinsicType::kCountOneBits, ParamType::kU32,
"metal::popcount"},
- IntrinsicData{semantic::Intrinsic::kCross, ParamType::kF32,
- "metal::cross"},
- IntrinsicData{semantic::Intrinsic::kDeterminant, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kCross, ParamType::kF32, "metal::cross"},
+ IntrinsicData{IntrinsicType::kDeterminant, ParamType::kF32,
"metal::determinant"},
- IntrinsicData{semantic::Intrinsic::kDistance, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kDistance, ParamType::kF32,
"metal::distance"},
- IntrinsicData{semantic::Intrinsic::kDot, ParamType::kF32, "metal::dot"},
- IntrinsicData{semantic::Intrinsic::kDpdx, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kDot, ParamType::kF32, "metal::dot"},
+ IntrinsicData{IntrinsicType::kDpdx, ParamType::kF32, "metal::dfdx"},
+ IntrinsicData{IntrinsicType::kDpdxCoarse, ParamType::kF32,
"metal::dfdx"},
- IntrinsicData{semantic::Intrinsic::kDpdxCoarse, ParamType::kF32,
- "metal::dfdx"},
- IntrinsicData{semantic::Intrinsic::kDpdxFine, ParamType::kF32,
- "metal::dfdx"},
- IntrinsicData{semantic::Intrinsic::kDpdy, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kDpdxFine, ParamType::kF32, "metal::dfdx"},
+ IntrinsicData{IntrinsicType::kDpdy, ParamType::kF32, "metal::dfdy"},
+ IntrinsicData{IntrinsicType::kDpdyCoarse, ParamType::kF32,
"metal::dfdy"},
- IntrinsicData{semantic::Intrinsic::kDpdyCoarse, ParamType::kF32,
- "metal::dfdy"},
- 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,
+ IntrinsicData{IntrinsicType::kDpdyFine, ParamType::kF32, "metal::dfdy"},
+ IntrinsicData{IntrinsicType::kExp, ParamType::kF32, "metal::exp"},
+ IntrinsicData{IntrinsicType::kExp2, ParamType::kF32, "metal::exp2"},
+ IntrinsicData{IntrinsicType::kFaceForward, ParamType::kF32,
"metal::faceforward"},
- 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,
+ IntrinsicData{IntrinsicType::kFloor, ParamType::kF32, "metal::floor"},
+ IntrinsicData{IntrinsicType::kFma, ParamType::kF32, "metal::fma"},
+ IntrinsicData{IntrinsicType::kFract, ParamType::kF32, "metal::fract"},
+ IntrinsicData{IntrinsicType::kFwidth, ParamType::kF32, "metal::fwidth"},
+ IntrinsicData{IntrinsicType::kFwidthCoarse, ParamType::kF32,
"metal::fwidth"},
- IntrinsicData{semantic::Intrinsic::kFwidthCoarse, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kFwidthFine, ParamType::kF32,
"metal::fwidth"},
- IntrinsicData{semantic::Intrinsic::kFwidthFine, ParamType::kF32,
- "metal::fwidth"},
- IntrinsicData{semantic::Intrinsic::kInverseSqrt, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kInverseSqrt, ParamType::kF32,
"metal::rsqrt"},
- IntrinsicData{semantic::Intrinsic::kIsFinite, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kIsFinite, ParamType::kF32,
"metal::isfinite"},
- IntrinsicData{semantic::Intrinsic::kIsInf, ParamType::kF32,
- "metal::isinf"},
- IntrinsicData{semantic::Intrinsic::kIsNan, ParamType::kF32,
- "metal::isnan"},
- IntrinsicData{semantic::Intrinsic::kIsNormal, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kIsInf, ParamType::kF32, "metal::isinf"},
+ IntrinsicData{IntrinsicType::kIsNan, ParamType::kF32, "metal::isnan"},
+ IntrinsicData{IntrinsicType::kIsNormal, ParamType::kF32,
"metal::isnormal"},
- IntrinsicData{semantic::Intrinsic::kLdexp, ParamType::kF32,
- "metal::ldexp"},
- IntrinsicData{semantic::Intrinsic::kLength, ParamType::kF32,
- "metal::length"},
- 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,
+ IntrinsicData{IntrinsicType::kLdexp, ParamType::kF32, "metal::ldexp"},
+ IntrinsicData{IntrinsicType::kLength, ParamType::kF32, "metal::length"},
+ IntrinsicData{IntrinsicType::kLog, ParamType::kF32, "metal::log"},
+ IntrinsicData{IntrinsicType::kLog2, ParamType::kF32, "metal::log2"},
+ IntrinsicData{IntrinsicType::kMax, ParamType::kF32, "metal::fmax"},
+ IntrinsicData{IntrinsicType::kMax, ParamType::kU32, "metal::max"},
+ IntrinsicData{IntrinsicType::kMin, ParamType::kF32, "metal::fmin"},
+ IntrinsicData{IntrinsicType::kMin, ParamType::kU32, "metal::min"},
+ IntrinsicData{IntrinsicType::kNormalize, ParamType::kF32,
"metal::normalize"},
- IntrinsicData{semantic::Intrinsic::kPack4x8Snorm, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kPack4x8Snorm, ParamType::kF32,
"metal::pack_float_to_snorm4x8"},
- IntrinsicData{semantic::Intrinsic::kPack4x8Unorm, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kPack4x8Unorm, ParamType::kF32,
"metal::pack_float_to_unorm4x8"},
- IntrinsicData{semantic::Intrinsic::kPack2x16Snorm, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kPack2x16Snorm, ParamType::kF32,
"metal::pack_float_to_snorm2x16"},
- IntrinsicData{semantic::Intrinsic::kPack2x16Unorm, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kPack2x16Unorm, ParamType::kF32,
"metal::pack_float_to_unorm2x16"},
- IntrinsicData{semantic::Intrinsic::kPow, ParamType::kF32, "metal::pow"},
- IntrinsicData{semantic::Intrinsic::kReflect, ParamType::kF32,
+ IntrinsicData{IntrinsicType::kPow, ParamType::kF32, "metal::pow"},
+ IntrinsicData{IntrinsicType::kReflect, ParamType::kF32,
"metal::reflect"},
- IntrinsicData{semantic::Intrinsic::kReverseBits, ParamType::kU32,
+ IntrinsicData{IntrinsicType::kReverseBits, ParamType::kU32,
"metal::reverse_bits"},
- IntrinsicData{semantic::Intrinsic::kRound, ParamType::kU32,
- "metal::round"},
- IntrinsicData{semantic::Intrinsic::kSelect, ParamType::kF32,
- "metal::select"},
- 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,
+ IntrinsicData{IntrinsicType::kRound, ParamType::kU32, "metal::round"},
+ IntrinsicData{IntrinsicType::kSelect, ParamType::kF32, "metal::select"},
+ IntrinsicData{IntrinsicType::kSign, ParamType::kF32, "metal::sign"},
+ IntrinsicData{IntrinsicType::kSin, ParamType::kF32, "metal::sin"},
+ IntrinsicData{IntrinsicType::kSinh, ParamType::kF32, "metal::sinh"},
+ IntrinsicData{IntrinsicType::kSmoothStep, ParamType::kF32,
"metal::smoothstep"},
- 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"}));
+ IntrinsicData{IntrinsicType::kSqrt, ParamType::kF32, "metal::sqrt"},
+ IntrinsicData{IntrinsicType::kStep, ParamType::kF32, "metal::step"},
+ IntrinsicData{IntrinsicType::kTan, ParamType::kF32, "metal::tan"},
+ IntrinsicData{IntrinsicType::kTanh, ParamType::kF32, "metal::tanh"},
+ IntrinsicData{IntrinsicType::kTrunc, ParamType::kF32, "metal::trunc"}));
TEST_F(MslGeneratorImplTest, Intrinsic_Call) {
Global("param1", ast::StorageClass::kFunction, ty.vec2<f32>());
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index a351357..ce32ebb 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -87,6 +87,8 @@
namespace spirv {
namespace {
+using IntrinsicType = semantic::IntrinsicType;
+
const char kGLSLstd450[] = "GLSL.std.450";
uint32_t size_of(const InstructionList& instructions) {
@@ -166,26 +168,25 @@
return type->As<type::Matrix>();
}
-uint32_t intrinsic_to_glsl_method(type::Type* type,
- semantic::Intrinsic intrinsic) {
+uint32_t intrinsic_to_glsl_method(type::Type* type, IntrinsicType intrinsic) {
switch (intrinsic) {
- case semantic::Intrinsic::kAbs:
+ case IntrinsicType::kAbs:
if (type->is_float_scalar_or_vector()) {
return GLSLstd450FAbs;
} else {
return GLSLstd450SAbs;
}
- case semantic::Intrinsic::kAcos:
+ case IntrinsicType::kAcos:
return GLSLstd450Acos;
- case semantic::Intrinsic::kAsin:
+ case IntrinsicType::kAsin:
return GLSLstd450Asin;
- case semantic::Intrinsic::kAtan:
+ case IntrinsicType::kAtan:
return GLSLstd450Atan;
- case semantic::Intrinsic::kAtan2:
+ case IntrinsicType::kAtan2:
return GLSLstd450Atan2;
- case semantic::Intrinsic::kCeil:
+ case IntrinsicType::kCeil:
return GLSLstd450Ceil;
- case semantic::Intrinsic::kClamp:
+ case IntrinsicType::kClamp:
if (type->is_float_scalar_or_vector()) {
return GLSLstd450NClamp;
} else if (type->is_unsigned_scalar_or_vector()) {
@@ -193,41 +194,41 @@
} else {
return GLSLstd450SClamp;
}
- case semantic::Intrinsic::kCos:
+ case IntrinsicType::kCos:
return GLSLstd450Cos;
- case semantic::Intrinsic::kCosh:
+ case IntrinsicType::kCosh:
return GLSLstd450Cosh;
- case semantic::Intrinsic::kCross:
+ case IntrinsicType::kCross:
return GLSLstd450Cross;
- case semantic::Intrinsic::kDeterminant:
+ case IntrinsicType::kDeterminant:
return GLSLstd450Determinant;
- case semantic::Intrinsic::kDistance:
+ case IntrinsicType::kDistance:
return GLSLstd450Distance;
- case semantic::Intrinsic::kExp:
+ case IntrinsicType::kExp:
return GLSLstd450Exp;
- case semantic::Intrinsic::kExp2:
+ case IntrinsicType::kExp2:
return GLSLstd450Exp2;
- case semantic::Intrinsic::kFaceForward:
+ case IntrinsicType::kFaceForward:
return GLSLstd450FaceForward;
- case semantic::Intrinsic::kFloor:
+ case IntrinsicType::kFloor:
return GLSLstd450Floor;
- case semantic::Intrinsic::kFma:
+ case IntrinsicType::kFma:
return GLSLstd450Fma;
- case semantic::Intrinsic::kFract:
+ case IntrinsicType::kFract:
return GLSLstd450Fract;
- case semantic::Intrinsic::kFrexp:
+ case IntrinsicType::kFrexp:
return GLSLstd450Frexp;
- case semantic::Intrinsic::kInverseSqrt:
+ case IntrinsicType::kInverseSqrt:
return GLSLstd450InverseSqrt;
- case semantic::Intrinsic::kLdexp:
+ case IntrinsicType::kLdexp:
return GLSLstd450Ldexp;
- case semantic::Intrinsic::kLength:
+ case IntrinsicType::kLength:
return GLSLstd450Length;
- case semantic::Intrinsic::kLog:
+ case IntrinsicType::kLog:
return GLSLstd450Log;
- case semantic::Intrinsic::kLog2:
+ case IntrinsicType::kLog2:
return GLSLstd450Log2;
- case semantic::Intrinsic::kMax:
+ case IntrinsicType::kMax:
if (type->is_float_scalar_or_vector()) {
return GLSLstd450NMax;
} else if (type->is_unsigned_scalar_or_vector()) {
@@ -235,7 +236,7 @@
} else {
return GLSLstd450SMax;
}
- case semantic::Intrinsic::kMin:
+ case IntrinsicType::kMin:
if (type->is_float_scalar_or_vector()) {
return GLSLstd450NMin;
} else if (type->is_unsigned_scalar_or_vector()) {
@@ -243,45 +244,45 @@
} else {
return GLSLstd450SMin;
}
- case semantic::Intrinsic::kMix:
+ case IntrinsicType::kMix:
return GLSLstd450FMix;
- case semantic::Intrinsic::kModf:
+ case IntrinsicType::kModf:
return GLSLstd450Modf;
- case semantic::Intrinsic::kNormalize:
+ case IntrinsicType::kNormalize:
return GLSLstd450Normalize;
- case semantic::Intrinsic::kPack4x8Snorm:
+ case IntrinsicType::kPack4x8Snorm:
return GLSLstd450PackSnorm4x8;
- case semantic::Intrinsic::kPack4x8Unorm:
+ case IntrinsicType::kPack4x8Unorm:
return GLSLstd450PackUnorm4x8;
- case semantic::Intrinsic::kPack2x16Snorm:
+ case IntrinsicType::kPack2x16Snorm:
return GLSLstd450PackSnorm2x16;
- case semantic::Intrinsic::kPack2x16Unorm:
+ case IntrinsicType::kPack2x16Unorm:
return GLSLstd450PackUnorm2x16;
- case semantic::Intrinsic::kPack2x16Float:
+ case IntrinsicType::kPack2x16Float:
return GLSLstd450PackHalf2x16;
- case semantic::Intrinsic::kPow:
+ case IntrinsicType::kPow:
return GLSLstd450Pow;
- case semantic::Intrinsic::kReflect:
+ case IntrinsicType::kReflect:
return GLSLstd450Reflect;
- case semantic::Intrinsic::kRound:
+ case IntrinsicType::kRound:
return GLSLstd450Round;
- case semantic::Intrinsic::kSign:
+ case IntrinsicType::kSign:
return GLSLstd450FSign;
- case semantic::Intrinsic::kSin:
+ case IntrinsicType::kSin:
return GLSLstd450Sin;
- case semantic::Intrinsic::kSinh:
+ case IntrinsicType::kSinh:
return GLSLstd450Sinh;
- case semantic::Intrinsic::kSmoothStep:
+ case IntrinsicType::kSmoothStep:
return GLSLstd450SmoothStep;
- case semantic::Intrinsic::kSqrt:
+ case IntrinsicType::kSqrt:
return GLSLstd450Sqrt;
- case semantic::Intrinsic::kStep:
+ case IntrinsicType::kStep:
return GLSLstd450Step;
- case semantic::Intrinsic::kTan:
+ case IntrinsicType::kTan:
return GLSLstd450Tan;
- case semantic::Intrinsic::kTanh:
+ case IntrinsicType::kTanh:
return GLSLstd450Tanh;
- case semantic::Intrinsic::kTrunc:
+ case IntrinsicType::kTrunc:
return GLSLstd450Trunc;
default:
break;
@@ -1897,11 +1898,11 @@
OperandList params = {Operand::Int(result_type_id), result};
spv::Op op = spv::Op::OpNop;
- if (intrinsic == semantic::Intrinsic::kAny) {
+ if (intrinsic == IntrinsicType::kAny) {
op = spv::Op::OpAny;
- } else if (intrinsic == semantic::Intrinsic::kAll) {
+ } else if (intrinsic == IntrinsicType::kAll) {
op = spv::Op::OpAll;
- } else if (intrinsic == semantic::Intrinsic::kArrayLength) {
+ } else if (intrinsic == IntrinsicType::kArrayLength) {
if (call->params().empty()) {
error_ = "missing param for runtime array length";
return 0;
@@ -1934,35 +1935,35 @@
return 0;
}
return result_id;
- } else if (intrinsic == semantic::Intrinsic::kCountOneBits) {
+ } else if (intrinsic == IntrinsicType::kCountOneBits) {
op = spv::Op::OpBitCount;
- } else if (intrinsic == semantic::Intrinsic::kDot) {
+ } else if (intrinsic == IntrinsicType::kDot) {
op = spv::Op::OpDot;
- } else if (intrinsic == semantic::Intrinsic::kDpdx) {
+ } else if (intrinsic == IntrinsicType::kDpdx) {
op = spv::Op::OpDPdx;
- } else if (intrinsic == semantic::Intrinsic::kDpdxCoarse) {
+ } else if (intrinsic == IntrinsicType::kDpdxCoarse) {
op = spv::Op::OpDPdxCoarse;
- } else if (intrinsic == semantic::Intrinsic::kDpdxFine) {
+ } else if (intrinsic == IntrinsicType::kDpdxFine) {
op = spv::Op::OpDPdxFine;
- } else if (intrinsic == semantic::Intrinsic::kDpdy) {
+ } else if (intrinsic == IntrinsicType::kDpdy) {
op = spv::Op::OpDPdy;
- } else if (intrinsic == semantic::Intrinsic::kDpdyCoarse) {
+ } else if (intrinsic == IntrinsicType::kDpdyCoarse) {
op = spv::Op::OpDPdyCoarse;
- } else if (intrinsic == semantic::Intrinsic::kDpdyFine) {
+ } else if (intrinsic == IntrinsicType::kDpdyFine) {
op = spv::Op::OpDPdyFine;
- } else if (intrinsic == semantic::Intrinsic::kFwidth) {
+ } else if (intrinsic == IntrinsicType::kFwidth) {
op = spv::Op::OpFwidth;
- } else if (intrinsic == semantic::Intrinsic::kFwidthCoarse) {
+ } else if (intrinsic == IntrinsicType::kFwidthCoarse) {
op = spv::Op::OpFwidthCoarse;
- } else if (intrinsic == semantic::Intrinsic::kFwidthFine) {
+ } else if (intrinsic == IntrinsicType::kFwidthFine) {
op = spv::Op::OpFwidthFine;
- } else if (intrinsic == semantic::Intrinsic::kIsInf) {
+ } else if (intrinsic == IntrinsicType::kIsInf) {
op = spv::Op::OpIsInf;
- } else if (intrinsic == semantic::Intrinsic::kIsNan) {
+ } else if (intrinsic == IntrinsicType::kIsNan) {
op = spv::Op::OpIsNan;
- } else if (intrinsic == semantic::Intrinsic::kReverseBits) {
+ } else if (intrinsic == IntrinsicType::kReverseBits) {
op = spv::Op::OpBitReverse;
- } else if (intrinsic == semantic::Intrinsic::kSelect) {
+ } else if (intrinsic == IntrinsicType::kSelect) {
op = spv::Op::OpSelect;
} else {
GenerateGLSLstd450Import();
@@ -2163,7 +2164,7 @@
};
switch (sem->intrinsic()) {
- case semantic::Intrinsic::kTextureDimensions: {
+ case IntrinsicType::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
@@ -2219,7 +2220,7 @@
}
break;
}
- case semantic::Intrinsic::kTextureNumLayers: {
+ case IntrinsicType::kTextureNumLayers: {
uint32_t spirv_dims = 0;
switch (texture_type->dim()) {
default:
@@ -2254,19 +2255,19 @@
}
break;
}
- case semantic::Intrinsic::kTextureNumLevels: {
+ case IntrinsicType::kTextureNumLevels: {
op = spv::Op::OpImageQueryLevels;
append_result_type_and_id_to_spirv_params();
spirv_params.emplace_back(gen_param(pidx.texture));
break;
}
- case semantic::Intrinsic::kTextureNumSamples: {
+ case IntrinsicType::kTextureNumSamples: {
op = spv::Op::OpImageQuerySamples;
append_result_type_and_id_to_spirv_params();
spirv_params.emplace_back(gen_param(pidx.texture));
break;
}
- case semantic::Intrinsic::kTextureLoad: {
+ case IntrinsicType::kTextureLoad: {
op = texture_type->Is<type::StorageTexture>() ? spv::Op::OpImageRead
: spv::Op::OpImageFetch;
append_result_type_and_id_to_spirv_params_for_read();
@@ -2287,7 +2288,7 @@
break;
}
- case semantic::Intrinsic::kTextureStore: {
+ case IntrinsicType::kTextureStore: {
op = spv::Op::OpImageWrite;
spirv_params.emplace_back(gen_param(pidx.texture));
if (!append_coords_to_spirv_params()) {
@@ -2296,7 +2297,7 @@
spirv_params.emplace_back(gen_param(pidx.value));
break;
}
- case semantic::Intrinsic::kTextureSample: {
+ case IntrinsicType::kTextureSample: {
op = spv::Op::OpImageSampleImplicitLod;
append_result_type_and_id_to_spirv_params_for_read();
if (!append_image_and_coords_to_spirv_params()) {
@@ -2304,7 +2305,7 @@
}
break;
}
- case semantic::Intrinsic::kTextureSampleBias: {
+ case IntrinsicType::kTextureSampleBias: {
op = spv::Op::OpImageSampleImplicitLod;
append_result_type_and_id_to_spirv_params_for_read();
if (!append_image_and_coords_to_spirv_params()) {
@@ -2315,7 +2316,7 @@
ImageOperand{SpvImageOperandsBiasMask, gen_param(pidx.bias)});
break;
}
- case semantic::Intrinsic::kTextureSampleLevel: {
+ case IntrinsicType::kTextureSampleLevel: {
op = spv::Op::OpImageSampleExplicitLod;
append_result_type_and_id_to_spirv_params_for_read();
if (!append_image_and_coords_to_spirv_params()) {
@@ -2339,7 +2340,7 @@
image_operands.emplace_back(ImageOperand{SpvImageOperandsLodMask, level});
break;
}
- case semantic::Intrinsic::kTextureSampleGrad: {
+ case IntrinsicType::kTextureSampleGrad: {
op = spv::Op::OpImageSampleExplicitLod;
append_result_type_and_id_to_spirv_params_for_read();
if (!append_image_and_coords_to_spirv_params()) {
@@ -2353,7 +2354,7 @@
ImageOperand{SpvImageOperandsGradMask, gen_param(pidx.ddy)});
break;
}
- case semantic::Intrinsic::kTextureSampleCompare: {
+ case IntrinsicType::kTextureSampleCompare: {
op = spv::Op::OpImageSampleDrefExplicitLod;
append_result_type_and_id_to_spirv_params();
if (!append_image_and_coords_to_spirv_params()) {
