encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

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>
This commit is contained in:
Ben Clayton 2021-02-08 19:53:42 +00:00 committed by Commit Bot service account
parent 2ddb1783c5
commit 052ab89a1e
14 changed files with 997 additions and 1066 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
src/reader/spirv/function.cc
View File

@ -470,19 +470,19 @@ std::string GetGlslStd450FuncName(uint32_t ext_opcode) {
} }
// Returns the WGSL standard library function intrinsic for the // Returns the WGSL standard library function intrinsic for the
// given instruction, or semantic::Intrinsic::kNone // given instruction, or semantic::IntrinsicType::kNone
semantic::Intrinsic GetIntrinsic(SpvOp opcode) { semantic::IntrinsicType GetIntrinsic(SpvOp opcode) {
switch (opcode) { switch (opcode) {
case SpvOpBitCount: case SpvOpBitCount:
return semantic::Intrinsic::kCountOneBits; return semantic::IntrinsicType::kCountOneBits;
case SpvOpBitReverse: case SpvOpBitReverse:
return semantic::Intrinsic::kReverseBits; return semantic::IntrinsicType::kReverseBits;
case SpvOpDot: case SpvOpDot:
return semantic::Intrinsic::kDot; return semantic::IntrinsicType::kDot;
default: default:
break; break;
} }
return semantic::Intrinsic::kNone; return semantic::IntrinsicType::kNone;
} }
// @param opcode a SPIR-V opcode // @param opcode a SPIR-V opcode
@ -3212,7 +3212,7 @@ TypedExpression FunctionEmitter::MaybeEmitCombinatorialValue(
} }
const auto intrinsic = GetIntrinsic(opcode); const auto intrinsic = GetIntrinsic(opcode);
if (intrinsic != semantic::Intrinsic::kNone) { if (intrinsic != semantic::IntrinsicType::kNone) {
return MakeIntrinsicCall(inst); return MakeIntrinsicCall(inst);
} }

8
src/semantic/call.h
View File

@ -40,16 +40,16 @@ class IntrinsicCall : public Castable<IntrinsicCall, Call> {
/// Constructor /// Constructor
/// @param return_type the return type of the call /// @param return_type the return type of the call
/// @param intrinsic the call target intrinsic /// @param intrinsic the call target intrinsic
IntrinsicCall(type::Type* return_type, Intrinsic intrinsic); IntrinsicCall(type::Type* return_type, IntrinsicType intrinsic);
/// Destructor /// Destructor
~IntrinsicCall() override; ~IntrinsicCall() override;
/// @returns the target intrinsic for the call /// @returns the target intrinsic for the call
Intrinsic intrinsic() const { return intrinsic_; } IntrinsicType intrinsic() const { return intrinsic_; }
private: private:
Intrinsic const intrinsic_; IntrinsicType const intrinsic_;
}; };
/// TextureIntrinsicCall holds semantic information for ast::CallExpression /// TextureIntrinsicCall holds semantic information for ast::CallExpression
@ -105,7 +105,7 @@ class TextureIntrinsicCall
/// @param intrinsic the call target intrinsic /// @param intrinsic the call target intrinsic
/// @param params the overload parameter info /// @param params the overload parameter info
TextureIntrinsicCall(type::Type* return_type, TextureIntrinsicCall(type::Type* return_type,
Intrinsic intrinsic, IntrinsicType intrinsic,
const Parameters& params); const Parameters& params);
/// Destructor /// Destructor

219
src/semantic/intrinsic.cc
View File

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

20
src/semantic/intrinsic.h
View File

@ -20,7 +20,7 @@
namespace tint { namespace tint {
namespace semantic { namespace semantic {
enum class Intrinsic { enum class IntrinsicType {
kNone = -1, kNone = -1,
kAbs, kAbs,
@ -104,48 +104,48 @@ enum class Intrinsic {
/// Emits the name of the intrinsic function. The spelling, /// Emits the name of the intrinsic function. The spelling,
/// including case, matches the name in the WGSL spec. /// 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 { namespace intrinsic {
/// Determines if the given `i` is a coarse derivative /// Determines if the given `i` is a coarse derivative
/// @param i the intrinsic /// @param i the intrinsic
/// @returns true if the given derivative is coarse. /// @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 /// Determines if the given `i` is a fine derivative
/// @param i the intrinsic /// @param i the intrinsic
/// @returns true if the given derivative is fine. /// @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 /// Determine if the given `i` is a derivative intrinsic
/// @param i the intrinsic /// @param i the intrinsic
/// @returns true if the given `i` is a derivative 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 /// Determines if the given `i` is a float classification intrinsic
/// @param i the intrinsic /// @param i the intrinsic
/// @returns true if the given `i` is a float 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 /// Determines if the given `i` is a texture operation intrinsic
/// @param i the intrinsic /// @param i the intrinsic
/// @returns true if the given `i` is a texture operation 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 /// Determines if the given `i` is a image query intrinsic
/// @param i the intrinsic /// @param i the intrinsic
/// @returns true if the given `i` is a image query 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 /// Determines if the given `i` is a data packing intrinsic
/// @param i the intrinsic /// @param i the intrinsic
/// @returns true if the given `i` is a data packing 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, /// @returns the name of the intrinsic function. The spelling, including case,
/// matches the name in the WGSL spec. /// matches the name in the WGSL spec.
const char* str(Intrinsic i); const char* str(IntrinsicType i);
} // namespace intrinsic } // namespace intrinsic
} // namespace semantic } // namespace semantic

4
src/semantic/sem_call.cc
View File

@ -25,13 +25,13 @@ Call::Call(type::Type* return_type) : Base(return_type) {}
Call::~Call() = default; Call::~Call() = default;
IntrinsicCall::IntrinsicCall(type::Type* return_type, Intrinsic intrinsic) IntrinsicCall::IntrinsicCall(type::Type* return_type, IntrinsicType intrinsic)
: Base(return_type), intrinsic_(intrinsic) {} : Base(return_type), intrinsic_(intrinsic) {}
IntrinsicCall::~IntrinsicCall() = default; IntrinsicCall::~IntrinsicCall() = default;
TextureIntrinsicCall::TextureIntrinsicCall(type::Type* return_type, TextureIntrinsicCall::TextureIntrinsicCall(type::Type* return_type,
Intrinsic intrinsic, IntrinsicType intrinsic,
const Parameters& params) const Parameters& params)
: Base(return_type, intrinsic), params_(params) {} : Base(return_type, intrinsic), params_(params) {}

333
src/type_determiner.cc
View File

@ -65,6 +65,11 @@
#include "src/type/void_type.h" #include "src/type/void_type.h"
namespace tint { namespace tint {
namespace {
using IntrinsicType = tint::semantic::IntrinsicType;
} // namespace
TypeDeterminer::TypeDeterminer(ProgramBuilder* builder) : builder_(builder) {} TypeDeterminer::TypeDeterminer(ProgramBuilder* builder) : builder_(builder) {}
@ -415,7 +420,7 @@ bool TypeDeterminer::DetermineCall(ast::CallExpression* call) {
auto name = builder_->Symbols().NameFor(ident->symbol()); auto name = builder_->Symbols().NameFor(ident->symbol());
auto intrinsic = MatchIntrinsic(name); auto intrinsic = MatchIntrinsic(name);
if (intrinsic != semantic::Intrinsic::kNone) { if (intrinsic != IntrinsicType::kNone) {
if (!DetermineIntrinsicCall(call, intrinsic)) { if (!DetermineIntrinsicCall(call, intrinsic)) {
return false; return false;
} }
@ -464,7 +469,7 @@ enum class IntrinsicDataType {
}; };
struct IntrinsicData { struct IntrinsicData {
semantic::Intrinsic intrinsic; IntrinsicType intrinsic;
IntrinsicDataType result_type; IntrinsicDataType result_type;
uint8_t result_vector_width; uint8_t result_vector_width;
uint8_t param_for_result_type; uint8_t param_for_result_type;
@ -473,74 +478,68 @@ struct IntrinsicData {
// Note, this isn't all the intrinsics. Some are handled specially before // Note, this isn't all the intrinsics. Some are handled specially before
// we get to the generic code. See the DetermineIntrinsic code below. // we get to the generic code. See the DetermineIntrinsic code below.
constexpr const IntrinsicData kIntrinsicData[] = { constexpr const IntrinsicData kIntrinsicData[] = {
{semantic::Intrinsic::kAbs, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kAbs, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kAcos, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kAcos, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kAll, IntrinsicDataType::kBool, 1, 0}, {IntrinsicType::kAll, IntrinsicDataType::kBool, 1, 0},
{semantic::Intrinsic::kAny, IntrinsicDataType::kBool, 1, 0}, {IntrinsicType::kAny, IntrinsicDataType::kBool, 1, 0},
{semantic::Intrinsic::kArrayLength, IntrinsicDataType::kUnsignedInteger, 1, {IntrinsicType::kArrayLength, IntrinsicDataType::kUnsignedInteger, 1, 0},
0}, {IntrinsicType::kAsin, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kAsin, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kAtan, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kAtan, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kAtan2, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kAtan2, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kCeil, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kCeil, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kClamp, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kClamp, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kCos, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kCos, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kCosh, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kCosh, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kCountOneBits, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kCountOneBits, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kCross, IntrinsicDataType::kFloat, 3, 0},
{semantic::Intrinsic::kCross, IntrinsicDataType::kFloat, 3, 0}, {IntrinsicType::kDeterminant, IntrinsicDataType::kFloat, 1, 0},
{semantic::Intrinsic::kDeterminant, IntrinsicDataType::kFloat, 1, 0}, {IntrinsicType::kDistance, IntrinsicDataType::kFloat, 1, 0},
{semantic::Intrinsic::kDistance, IntrinsicDataType::kFloat, 1, 0}, {IntrinsicType::kDpdx, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kDpdx, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kDpdxCoarse, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kDpdxCoarse, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kDpdxFine, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kDpdxFine, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kDpdy, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kDpdy, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kDpdyCoarse, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kDpdyCoarse, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kDpdyFine, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kDpdyFine, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kDot, IntrinsicDataType::kFloat, 1, 0},
{semantic::Intrinsic::kDot, IntrinsicDataType::kFloat, 1, 0}, {IntrinsicType::kExp, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kExp, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kExp2, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kExp2, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kFaceForward, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kFaceForward, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kFloor, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kFloor, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kFwidth, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kFwidth, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kFwidthCoarse, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kFwidthCoarse, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kFwidthFine, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kFwidthFine, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kFma, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kFma, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kFract, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kFract, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kFrexp, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kFrexp, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kInverseSqrt, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kInverseSqrt, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kLdexp, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kLdexp, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kLength, IntrinsicDataType::kFloat, 1, 0},
{semantic::Intrinsic::kLength, IntrinsicDataType::kFloat, 1, 0}, {IntrinsicType::kLog, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kLog, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kLog2, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kLog2, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kMax, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kMax, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kMin, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kMin, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kMix, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kMix, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kModf, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kModf, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kNormalize, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kNormalize, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kPack4x8Snorm, IntrinsicDataType::kUnsignedInteger, 1, 0},
{semantic::Intrinsic::kPack4x8Snorm, IntrinsicDataType::kUnsignedInteger, 1, {IntrinsicType::kPack4x8Unorm, IntrinsicDataType::kUnsignedInteger, 1, 0},
0}, {IntrinsicType::kPack2x16Snorm, IntrinsicDataType::kUnsignedInteger, 1, 0},
{semantic::Intrinsic::kPack4x8Unorm, IntrinsicDataType::kUnsignedInteger, 1, {IntrinsicType::kPack2x16Unorm, IntrinsicDataType::kUnsignedInteger, 1, 0},
0}, {IntrinsicType::kPack2x16Float, IntrinsicDataType::kUnsignedInteger, 1, 0},
{semantic::Intrinsic::kPack2x16Snorm, IntrinsicDataType::kUnsignedInteger, {IntrinsicType::kPow, IntrinsicDataType::kDependent, 0, 0},
1, 0}, {IntrinsicType::kReflect, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kPack2x16Unorm, IntrinsicDataType::kUnsignedInteger, {IntrinsicType::kReverseBits, IntrinsicDataType::kDependent, 0, 0},
1, 0}, {IntrinsicType::kRound, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kPack2x16Float, IntrinsicDataType::kUnsignedInteger, {IntrinsicType::kSelect, IntrinsicDataType::kDependent, 0, 0},
1, 0}, {IntrinsicType::kSign, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kPow, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kSin, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kReflect, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kSinh, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kReverseBits, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kSmoothStep, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kRound, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kSqrt, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kSelect, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kStep, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kSign, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kTan, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kSin, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kTanh, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kSinh, IntrinsicDataType::kDependent, 0, 0}, {IntrinsicType::kTrunc, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kSmoothStep, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kSqrt, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kStep, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kTan, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kTanh, IntrinsicDataType::kDependent, 0, 0},
{semantic::Intrinsic::kTrunc, IntrinsicDataType::kDependent, 0, 0},
}; };
constexpr const uint32_t kIntrinsicDataCount = constexpr const uint32_t kIntrinsicDataCount =
@ -549,7 +548,7 @@ constexpr const uint32_t kIntrinsicDataCount =
} // namespace } // namespace
bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call, bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
semantic::Intrinsic intrinsic) { IntrinsicType intrinsic) {
auto create_sem = [&](type::Type* result) { auto create_sem = [&](type::Type* result) {
auto* sem = builder_->create<semantic::IntrinsicCall>(result, intrinsic); auto* sem = builder_->create<semantic::IntrinsicCall>(result, intrinsic);
builder_->Sem().Add(call, sem); builder_->Sem().Add(call, sem);
@ -586,18 +585,18 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
bool is_array = type::IsTextureArray(texture->dim()); bool is_array = type::IsTextureArray(texture->dim());
bool is_multisampled = texture->Is<type::MultisampledTexture>(); bool is_multisampled = texture->Is<type::MultisampledTexture>();
switch (intrinsic) { switch (intrinsic) {
case semantic::Intrinsic::kTextureDimensions: case IntrinsicType::kTextureDimensions:
param.idx.texture = param.count++; param.idx.texture = param.count++;
if (call->params().size() > param.count) { if (call->params().size() > param.count) {
param.idx.level = param.count++; param.idx.level = param.count++;
} }
break; break;
case semantic::Intrinsic::kTextureNumLayers: case IntrinsicType::kTextureNumLayers:
case semantic::Intrinsic::kTextureNumLevels: case IntrinsicType::kTextureNumLevels:
case semantic::Intrinsic::kTextureNumSamples: case IntrinsicType::kTextureNumSamples:
param.idx.texture = param.count++; param.idx.texture = param.count++;
break; break;
case semantic::Intrinsic::kTextureLoad: case IntrinsicType::kTextureLoad:
param.idx.texture = param.count++; param.idx.texture = param.count++;
param.idx.coords = param.count++; param.idx.coords = param.count++;
if (is_array) { if (is_array) {
@ -611,7 +610,7 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
} }
} }
break; break;
case semantic::Intrinsic::kTextureSample: case IntrinsicType::kTextureSample:
param.idx.texture = param.count++; param.idx.texture = param.count++;
param.idx.sampler = param.count++; param.idx.sampler = param.count++;
param.idx.coords = param.count++; param.idx.coords = param.count++;
@ -622,7 +621,7 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
param.idx.offset = param.count++; param.idx.offset = param.count++;
} }
break; break;
case semantic::Intrinsic::kTextureSampleBias: case IntrinsicType::kTextureSampleBias:
param.idx.texture = param.count++; param.idx.texture = param.count++;
param.idx.sampler = param.count++; param.idx.sampler = param.count++;
param.idx.coords = param.count++; param.idx.coords = param.count++;
@ -634,7 +633,7 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
param.idx.offset = param.count++; param.idx.offset = param.count++;
} }
break; break;
case semantic::Intrinsic::kTextureSampleLevel: case IntrinsicType::kTextureSampleLevel:
param.idx.texture = param.count++; param.idx.texture = param.count++;
param.idx.sampler = param.count++; param.idx.sampler = param.count++;
param.idx.coords = param.count++; param.idx.coords = param.count++;
@ -646,7 +645,7 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
param.idx.offset = param.count++; param.idx.offset = param.count++;
} }
break; break;
case semantic::Intrinsic::kTextureSampleCompare: case IntrinsicType::kTextureSampleCompare:
param.idx.texture = param.count++; param.idx.texture = param.count++;
param.idx.sampler = param.count++; param.idx.sampler = param.count++;
param.idx.coords = param.count++; param.idx.coords = param.count++;
@ -658,7 +657,7 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
param.idx.offset = param.count++; param.idx.offset = param.count++;
} }
break; break;
case semantic::Intrinsic::kTextureSampleGrad: case IntrinsicType::kTextureSampleGrad:
param.idx.texture = param.count++; param.idx.texture = param.count++;
param.idx.sampler = param.count++; param.idx.sampler = param.count++;
param.idx.coords = param.count++; param.idx.coords = param.count++;
@ -671,7 +670,7 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
param.idx.offset = param.count++; param.idx.offset = param.count++;
} }
break; break;
case semantic::Intrinsic::kTextureStore: case IntrinsicType::kTextureStore:
param.idx.texture = param.count++; param.idx.texture = param.count++;
param.idx.coords = param.count++; param.idx.coords = param.count++;
if (is_array) { if (is_array) {
@ -697,7 +696,7 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
// Set the function return type // Set the function return type
type::Type* return_type = nullptr; type::Type* return_type = nullptr;
switch (intrinsic) { switch (intrinsic) {
case semantic::Intrinsic::kTextureDimensions: { case IntrinsicType::kTextureDimensions: {
auto* i32 = builder_->create<type::I32>(); auto* i32 = builder_->create<type::I32>();
switch (texture->dim()) { switch (texture->dim()) {
default: default:
@ -719,12 +718,12 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
} }
break; break;
} }
case semantic::Intrinsic::kTextureNumLayers: case IntrinsicType::kTextureNumLayers:
case semantic::Intrinsic::kTextureNumLevels: case IntrinsicType::kTextureNumLevels:
case semantic::Intrinsic::kTextureNumSamples: case IntrinsicType::kTextureNumSamples:
return_type = builder_->create<type::I32>(); return_type = builder_->create<type::I32>();
break; break;
case semantic::Intrinsic::kTextureStore: case IntrinsicType::kTextureStore:
return_type = builder_->create<type::Void>(); return_type = builder_->create<type::Void>();
break; break;
default: { default: {
@ -848,7 +847,7 @@ bool TypeDeterminer::DetermineIdentifier(ast::IdentifierExpression* expr) {
} }
std::string name = builder_->Symbols().NameFor(symbol); 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). // Identifier is to an intrinsic function, which has no type (currently).
return true; return true;
} }
@ -858,163 +857,163 @@ bool TypeDeterminer::DetermineIdentifier(ast::IdentifierExpression* expr) {
return false; return false;
} }
semantic::Intrinsic TypeDeterminer::MatchIntrinsic(const std::string& name) { IntrinsicType TypeDeterminer::MatchIntrinsic(const std::string& name) {
if (name == "abs") { if (name == "abs") {
return semantic::Intrinsic::kAbs; return IntrinsicType::kAbs;
} else if (name == "acos") { } else if (name == "acos") {
return semantic::Intrinsic::kAcos; return IntrinsicType::kAcos;
} else if (name == "all") { } else if (name == "all") {
return semantic::Intrinsic::kAll; return IntrinsicType::kAll;
} else if (name == "any") { } else if (name == "any") {
return semantic::Intrinsic::kAny; return IntrinsicType::kAny;
} else if (name == "arrayLength") { } else if (name == "arrayLength") {
return semantic::Intrinsic::kArrayLength; return IntrinsicType::kArrayLength;
} else if (name == "asin") { } else if (name == "asin") {
return semantic::Intrinsic::kAsin; return IntrinsicType::kAsin;
} else if (name == "atan") { } else if (name == "atan") {
return semantic::Intrinsic::kAtan; return IntrinsicType::kAtan;
} else if (name == "atan2") { } else if (name == "atan2") {
return semantic::Intrinsic::kAtan2; return IntrinsicType::kAtan2;
} else if (name == "ceil") { } else if (name == "ceil") {
return semantic::Intrinsic::kCeil; return IntrinsicType::kCeil;
} else if (name == "clamp") { } else if (name == "clamp") {
return semantic::Intrinsic::kClamp; return IntrinsicType::kClamp;
} else if (name == "cos") { } else if (name == "cos") {
return semantic::Intrinsic::kCos; return IntrinsicType::kCos;
} else if (name == "cosh") { } else if (name == "cosh") {
return semantic::Intrinsic::kCosh; return IntrinsicType::kCosh;
} else if (name == "countOneBits") { } else if (name == "countOneBits") {
return semantic::Intrinsic::kCountOneBits; return IntrinsicType::kCountOneBits;
} else if (name == "cross") { } else if (name == "cross") {
return semantic::Intrinsic::kCross; return IntrinsicType::kCross;
} else if (name == "determinant") { } else if (name == "determinant") {
return semantic::Intrinsic::kDeterminant; return IntrinsicType::kDeterminant;
} else if (name == "distance") { } else if (name == "distance") {
return semantic::Intrinsic::kDistance; return IntrinsicType::kDistance;
} else if (name == "dot") { } else if (name == "dot") {
return semantic::Intrinsic::kDot; return IntrinsicType::kDot;
} else if (name == "dpdx") { } else if (name == "dpdx") {
return semantic::Intrinsic::kDpdx; return IntrinsicType::kDpdx;
} else if (name == "dpdxCoarse") { } else if (name == "dpdxCoarse") {
return semantic::Intrinsic::kDpdxCoarse; return IntrinsicType::kDpdxCoarse;
} else if (name == "dpdxFine") { } else if (name == "dpdxFine") {
return semantic::Intrinsic::kDpdxFine; return IntrinsicType::kDpdxFine;
} else if (name == "dpdy") { } else if (name == "dpdy") {
return semantic::Intrinsic::kDpdy; return IntrinsicType::kDpdy;
} else if (name == "dpdyCoarse") { } else if (name == "dpdyCoarse") {
return semantic::Intrinsic::kDpdyCoarse; return IntrinsicType::kDpdyCoarse;
} else if (name == "dpdyFine") { } else if (name == "dpdyFine") {
return semantic::Intrinsic::kDpdyFine; return IntrinsicType::kDpdyFine;
} else if (name == "exp") { } else if (name == "exp") {
return semantic::Intrinsic::kExp; return IntrinsicType::kExp;
} else if (name == "exp2") { } else if (name == "exp2") {
return semantic::Intrinsic::kExp2; return IntrinsicType::kExp2;
} else if (name == "faceForward") { } else if (name == "faceForward") {
return semantic::Intrinsic::kFaceForward; return IntrinsicType::kFaceForward;
} else if (name == "floor") { } else if (name == "floor") {
return semantic::Intrinsic::kFloor; return IntrinsicType::kFloor;
} else if (name == "fma") { } else if (name == "fma") {
return semantic::Intrinsic::kFma; return IntrinsicType::kFma;
} else if (name == "fract") { } else if (name == "fract") {
return semantic::Intrinsic::kFract; return IntrinsicType::kFract;
} else if (name == "frexp") { } else if (name == "frexp") {
return semantic::Intrinsic::kFrexp; return IntrinsicType::kFrexp;
} else if (name == "fwidth") { } else if (name == "fwidth") {
return semantic::Intrinsic::kFwidth; return IntrinsicType::kFwidth;
} else if (name == "fwidthCoarse") { } else if (name == "fwidthCoarse") {
return semantic::Intrinsic::kFwidthCoarse; return IntrinsicType::kFwidthCoarse;
} else if (name == "fwidthFine") { } else if (name == "fwidthFine") {
return semantic::Intrinsic::kFwidthFine; return IntrinsicType::kFwidthFine;
} else if (name == "inverseSqrt") { } else if (name == "inverseSqrt") {
return semantic::Intrinsic::kInverseSqrt; return IntrinsicType::kInverseSqrt;
} else if (name == "isFinite") { } else if (name == "isFinite") {
return semantic::Intrinsic::kIsFinite; return IntrinsicType::kIsFinite;
} else if (name == "isInf") { } else if (name == "isInf") {
return semantic::Intrinsic::kIsInf; return IntrinsicType::kIsInf;
} else if (name == "isNan") { } else if (name == "isNan") {
return semantic::Intrinsic::kIsNan; return IntrinsicType::kIsNan;
} else if (name == "isNormal") { } else if (name == "isNormal") {
return semantic::Intrinsic::kIsNormal; return IntrinsicType::kIsNormal;
} else if (name == "ldexp") { } else if (name == "ldexp") {
return semantic::Intrinsic::kLdexp; return IntrinsicType::kLdexp;
} else if (name == "length") { } else if (name == "length") {
return semantic::Intrinsic::kLength; return IntrinsicType::kLength;
} else if (name == "log") { } else if (name == "log") {
return semantic::Intrinsic::kLog; return IntrinsicType::kLog;
} else if (name == "log2") { } else if (name == "log2") {
return semantic::Intrinsic::kLog2; return IntrinsicType::kLog2;
} else if (name == "max") { } else if (name == "max") {
return semantic::Intrinsic::kMax; return IntrinsicType::kMax;
} else if (name == "min") { } else if (name == "min") {
return semantic::Intrinsic::kMin; return IntrinsicType::kMin;
} else if (name == "mix") { } else if (name == "mix") {
return semantic::Intrinsic::kMix; return IntrinsicType::kMix;
} else if (name == "modf") { } else if (name == "modf") {
return semantic::Intrinsic::kModf; return IntrinsicType::kModf;
} else if (name == "normalize") { } else if (name == "normalize") {
return semantic::Intrinsic::kNormalize; return IntrinsicType::kNormalize;
} else if (name == "pack4x8snorm") { } else if (name == "pack4x8snorm") {
return semantic::Intrinsic::kPack4x8Snorm; return IntrinsicType::kPack4x8Snorm;
} else if (name == "pack4x8unorm") { } else if (name == "pack4x8unorm") {
return semantic::Intrinsic::kPack4x8Unorm; return IntrinsicType::kPack4x8Unorm;
} else if (name == "pack2x16snorm") { } else if (name == "pack2x16snorm") {
return semantic::Intrinsic::kPack2x16Snorm; return IntrinsicType::kPack2x16Snorm;
} else if (name == "pack2x16unorm") { } else if (name == "pack2x16unorm") {
return semantic::Intrinsic::kPack2x16Unorm; return IntrinsicType::kPack2x16Unorm;
} else if (name == "pack2x16float") { } else if (name == "pack2x16float") {
return semantic::Intrinsic::kPack2x16Float; return IntrinsicType::kPack2x16Float;
} else if (name == "pow") { } else if (name == "pow") {
return semantic::Intrinsic::kPow; return IntrinsicType::kPow;
} else if (name == "reflect") { } else if (name == "reflect") {
return semantic::Intrinsic::kReflect; return IntrinsicType::kReflect;
} else if (name == "reverseBits") { } else if (name == "reverseBits") {
return semantic::Intrinsic::kReverseBits; return IntrinsicType::kReverseBits;
} else if (name == "round") { } else if (name == "round") {
return semantic::Intrinsic::kRound; return IntrinsicType::kRound;
} else if (name == "select") { } else if (name == "select") {
return semantic::Intrinsic::kSelect; return IntrinsicType::kSelect;
} else if (name == "sign") { } else if (name == "sign") {
return semantic::Intrinsic::kSign; return IntrinsicType::kSign;
} else if (name == "sin") { } else if (name == "sin") {
return semantic::Intrinsic::kSin; return IntrinsicType::kSin;
} else if (name == "sinh") { } else if (name == "sinh") {
return semantic::Intrinsic::kSinh; return IntrinsicType::kSinh;
} else if (name == "smoothStep") { } else if (name == "smoothStep") {
return semantic::Intrinsic::kSmoothStep; return IntrinsicType::kSmoothStep;
} else if (name == "sqrt") { } else if (name == "sqrt") {
return semantic::Intrinsic::kSqrt; return IntrinsicType::kSqrt;
} else if (name == "step") { } else if (name == "step") {
return semantic::Intrinsic::kStep; return IntrinsicType::kStep;
} else if (name == "tan") { } else if (name == "tan") {
return semantic::Intrinsic::kTan; return IntrinsicType::kTan;
} else if (name == "tanh") { } else if (name == "tanh") {
return semantic::Intrinsic::kTanh; return IntrinsicType::kTanh;
} else if (name == "textureDimensions") { } else if (name == "textureDimensions") {
return semantic::Intrinsic::kTextureDimensions; return IntrinsicType::kTextureDimensions;
} else if (name == "textureNumLayers") { } else if (name == "textureNumLayers") {
return semantic::Intrinsic::kTextureNumLayers; return IntrinsicType::kTextureNumLayers;
} else if (name == "textureNumLevels") { } else if (name == "textureNumLevels") {
return semantic::Intrinsic::kTextureNumLevels; return IntrinsicType::kTextureNumLevels;
} else if (name == "textureNumSamples") { } else if (name == "textureNumSamples") {
return semantic::Intrinsic::kTextureNumSamples; return IntrinsicType::kTextureNumSamples;
} else if (name == "textureLoad") { } else if (name == "textureLoad") {
return semantic::Intrinsic::kTextureLoad; return IntrinsicType::kTextureLoad;
} else if (name == "textureStore") { } else if (name == "textureStore") {
return semantic::Intrinsic::kTextureStore; return IntrinsicType::kTextureStore;
} else if (name == "textureSample") { } else if (name == "textureSample") {
return semantic::Intrinsic::kTextureSample; return IntrinsicType::kTextureSample;
} else if (name == "textureSampleBias") { } else if (name == "textureSampleBias") {
return semantic::Intrinsic::kTextureSampleBias; return IntrinsicType::kTextureSampleBias;
} else if (name == "textureSampleCompare") { } else if (name == "textureSampleCompare") {
return semantic::Intrinsic::kTextureSampleCompare; return IntrinsicType::kTextureSampleCompare;
} else if (name == "textureSampleGrad") { } else if (name == "textureSampleGrad") {
return semantic::Intrinsic::kTextureSampleGrad; return IntrinsicType::kTextureSampleGrad;
} else if (name == "textureSampleLevel") { } else if (name == "textureSampleLevel") {
return semantic::Intrinsic::kTextureSampleLevel; return IntrinsicType::kTextureSampleLevel;
} else if (name == "trunc") { } else if (name == "trunc") {
return semantic::Intrinsic::kTrunc; return IntrinsicType::kTrunc;
} }
return semantic::Intrinsic::kNone; return IntrinsicType::kNone;
} }
bool TypeDeterminer::DetermineMemberAccessor( bool TypeDeterminer::DetermineMemberAccessor(

8
src/type_determiner.h
View File

@ -67,9 +67,9 @@ class TypeDeterminer {
bool Determine(); bool Determine();
/// @param name the function name to try and match as an intrinsic. /// @param name the function name to try and match as an intrinsic.
/// @return the semantic::Intrinsic for the given name. If `name` does not /// @return the semantic::IntrinsicType for the given name. If `name` does not
/// match an intrinsic, returns semantic::Intrinsic::kNone /// match an intrinsic, returns semantic::IntrinsicType::kNone
static semantic::Intrinsic MatchIntrinsic(const std::string& name); static semantic::IntrinsicType MatchIntrinsic(const std::string& name);
private: private:
template <typename T> template <typename T>
@ -177,7 +177,7 @@ class TypeDeterminer {
bool DetermineConstructor(ast::ConstructorExpression* expr); bool DetermineConstructor(ast::ConstructorExpression* expr);
bool DetermineIdentifier(ast::IdentifierExpression* expr); bool DetermineIdentifier(ast::IdentifierExpression* expr);
bool DetermineIntrinsicCall(ast::CallExpression* call, bool DetermineIntrinsicCall(ast::CallExpression* call,
semantic::Intrinsic intrinsic); semantic::IntrinsicType intrinsic);
bool DetermineMemberAccessor(ast::MemberAccessorExpression* expr); bool DetermineMemberAccessor(ast::MemberAccessorExpression* expr);
bool DetermineUnaryOp(ast::UnaryOpExpression* expr); bool DetermineUnaryOp(ast::UnaryOpExpression* expr);

247
src/type_determiner_test.cc
View File

@ -75,6 +75,8 @@
namespace tint { namespace tint {
namespace { namespace {
using IntrinsicType = semantic::IntrinsicType;
class FakeStmt : public ast::Statement { class FakeStmt : public ast::Statement {
public: public:
explicit FakeStmt(Source source) : ast::Statement(source) {} explicit FakeStmt(Source source) : ast::Statement(source) {}
@ -1618,7 +1620,7 @@ TEST_F(TypeDeterminerTest, StorageClass_NonFunctionClassError) {
struct IntrinsicData { struct IntrinsicData {
const char* name; const char* name;
semantic::Intrinsic intrinsic; IntrinsicType intrinsic;
}; };
inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) { inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
out << data.name; out << data.name;
@ -1634,89 +1636,82 @@ INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest, TypeDeterminerTest,
IntrinsicDataTest, IntrinsicDataTest,
testing::Values( testing::Values(
IntrinsicData{"abs", semantic::Intrinsic::kAbs}, IntrinsicData{"abs", IntrinsicType::kAbs},
IntrinsicData{"acos", semantic::Intrinsic::kAcos}, IntrinsicData{"acos", IntrinsicType::kAcos},
IntrinsicData{"all", semantic::Intrinsic::kAll}, IntrinsicData{"all", IntrinsicType::kAll},
IntrinsicData{"any", semantic::Intrinsic::kAny}, IntrinsicData{"any", IntrinsicType::kAny},
IntrinsicData{"arrayLength", semantic::Intrinsic::kArrayLength}, IntrinsicData{"arrayLength", IntrinsicType::kArrayLength},
IntrinsicData{"asin", semantic::Intrinsic::kAsin}, IntrinsicData{"asin", IntrinsicType::kAsin},
IntrinsicData{"atan", semantic::Intrinsic::kAtan}, IntrinsicData{"atan", IntrinsicType::kAtan},
IntrinsicData{"atan2", semantic::Intrinsic::kAtan2}, IntrinsicData{"atan2", IntrinsicType::kAtan2},
IntrinsicData{"ceil", semantic::Intrinsic::kCeil}, IntrinsicData{"ceil", IntrinsicType::kCeil},
IntrinsicData{"clamp", semantic::Intrinsic::kClamp}, IntrinsicData{"clamp", IntrinsicType::kClamp},
IntrinsicData{"cos", semantic::Intrinsic::kCos}, IntrinsicData{"cos", IntrinsicType::kCos},
IntrinsicData{"cosh", semantic::Intrinsic::kCosh}, IntrinsicData{"cosh", IntrinsicType::kCosh},
IntrinsicData{"countOneBits", semantic::Intrinsic::kCountOneBits}, IntrinsicData{"countOneBits", IntrinsicType::kCountOneBits},
IntrinsicData{"cross", semantic::Intrinsic::kCross}, IntrinsicData{"cross", IntrinsicType::kCross},
IntrinsicData{"determinant", semantic::Intrinsic::kDeterminant}, IntrinsicData{"determinant", IntrinsicType::kDeterminant},
IntrinsicData{"distance", semantic::Intrinsic::kDistance}, IntrinsicData{"distance", IntrinsicType::kDistance},
IntrinsicData{"dot", semantic::Intrinsic::kDot}, IntrinsicData{"dot", IntrinsicType::kDot},
IntrinsicData{"dpdx", semantic::Intrinsic::kDpdx}, IntrinsicData{"dpdx", IntrinsicType::kDpdx},
IntrinsicData{"dpdxCoarse", semantic::Intrinsic::kDpdxCoarse}, IntrinsicData{"dpdxCoarse", IntrinsicType::kDpdxCoarse},
IntrinsicData{"dpdxFine", semantic::Intrinsic::kDpdxFine}, IntrinsicData{"dpdxFine", IntrinsicType::kDpdxFine},
IntrinsicData{"dpdy", semantic::Intrinsic::kDpdy}, IntrinsicData{"dpdy", IntrinsicType::kDpdy},
IntrinsicData{"dpdyCoarse", semantic::Intrinsic::kDpdyCoarse}, IntrinsicData{"dpdyCoarse", IntrinsicType::kDpdyCoarse},
IntrinsicData{"dpdyFine", semantic::Intrinsic::kDpdyFine}, IntrinsicData{"dpdyFine", IntrinsicType::kDpdyFine},
IntrinsicData{"exp", semantic::Intrinsic::kExp}, IntrinsicData{"exp", IntrinsicType::kExp},
IntrinsicData{"exp2", semantic::Intrinsic::kExp2}, IntrinsicData{"exp2", IntrinsicType::kExp2},
IntrinsicData{"faceForward", semantic::Intrinsic::kFaceForward}, IntrinsicData{"faceForward", IntrinsicType::kFaceForward},
IntrinsicData{"floor", semantic::Intrinsic::kFloor}, IntrinsicData{"floor", IntrinsicType::kFloor},
IntrinsicData{"fma", semantic::Intrinsic::kFma}, IntrinsicData{"fma", IntrinsicType::kFma},
IntrinsicData{"fract", semantic::Intrinsic::kFract}, IntrinsicData{"fract", IntrinsicType::kFract},
IntrinsicData{"frexp", semantic::Intrinsic::kFrexp}, IntrinsicData{"frexp", IntrinsicType::kFrexp},
IntrinsicData{"fwidth", semantic::Intrinsic::kFwidth}, IntrinsicData{"fwidth", IntrinsicType::kFwidth},
IntrinsicData{"fwidthCoarse", semantic::Intrinsic::kFwidthCoarse}, IntrinsicData{"fwidthCoarse", IntrinsicType::kFwidthCoarse},
IntrinsicData{"fwidthFine", semantic::Intrinsic::kFwidthFine}, IntrinsicData{"fwidthFine", IntrinsicType::kFwidthFine},
IntrinsicData{"inverseSqrt", semantic::Intrinsic::kInverseSqrt}, IntrinsicData{"inverseSqrt", IntrinsicType::kInverseSqrt},
IntrinsicData{"isFinite", semantic::Intrinsic::kIsFinite}, IntrinsicData{"isFinite", IntrinsicType::kIsFinite},
IntrinsicData{"isInf", semantic::Intrinsic::kIsInf}, IntrinsicData{"isInf", IntrinsicType::kIsInf},
IntrinsicData{"isNan", semantic::Intrinsic::kIsNan}, IntrinsicData{"isNan", IntrinsicType::kIsNan},
IntrinsicData{"isNormal", semantic::Intrinsic::kIsNormal}, IntrinsicData{"isNormal", IntrinsicType::kIsNormal},
IntrinsicData{"ldexp", semantic::Intrinsic::kLdexp}, IntrinsicData{"ldexp", IntrinsicType::kLdexp},
IntrinsicData{"length", semantic::Intrinsic::kLength}, IntrinsicData{"length", IntrinsicType::kLength},
IntrinsicData{"log", semantic::Intrinsic::kLog}, IntrinsicData{"log", IntrinsicType::kLog},
IntrinsicData{"log2", semantic::Intrinsic::kLog2}, IntrinsicData{"log2", IntrinsicType::kLog2},
IntrinsicData{"max", semantic::Intrinsic::kMax}, IntrinsicData{"max", IntrinsicType::kMax},
IntrinsicData{"min", semantic::Intrinsic::kMin}, IntrinsicData{"min", IntrinsicType::kMin},
IntrinsicData{"mix", semantic::Intrinsic::kMix}, IntrinsicData{"mix", IntrinsicType::kMix},
IntrinsicData{"modf", semantic::Intrinsic::kModf}, IntrinsicData{"modf", IntrinsicType::kModf},
IntrinsicData{"normalize", semantic::Intrinsic::kNormalize}, IntrinsicData{"normalize", IntrinsicType::kNormalize},
IntrinsicData{"pow", semantic::Intrinsic::kPow}, IntrinsicData{"pow", IntrinsicType::kPow},
IntrinsicData{"reflect", semantic::Intrinsic::kReflect}, IntrinsicData{"reflect", IntrinsicType::kReflect},
IntrinsicData{"reverseBits", semantic::Intrinsic::kReverseBits}, IntrinsicData{"reverseBits", IntrinsicType::kReverseBits},
IntrinsicData{"round", semantic::Intrinsic::kRound}, IntrinsicData{"round", IntrinsicType::kRound},
IntrinsicData{"select", semantic::Intrinsic::kSelect}, IntrinsicData{"select", IntrinsicType::kSelect},
IntrinsicData{"sign", semantic::Intrinsic::kSign}, IntrinsicData{"sign", IntrinsicType::kSign},
IntrinsicData{"sin", semantic::Intrinsic::kSin}, IntrinsicData{"sin", IntrinsicType::kSin},
IntrinsicData{"sinh", semantic::Intrinsic::kSinh}, IntrinsicData{"sinh", IntrinsicType::kSinh},
IntrinsicData{"smoothStep", semantic::Intrinsic::kSmoothStep}, IntrinsicData{"smoothStep", IntrinsicType::kSmoothStep},
IntrinsicData{"sqrt", semantic::Intrinsic::kSqrt}, IntrinsicData{"sqrt", IntrinsicType::kSqrt},
IntrinsicData{"step", semantic::Intrinsic::kStep}, IntrinsicData{"step", IntrinsicType::kStep},
IntrinsicData{"tan", semantic::Intrinsic::kTan}, IntrinsicData{"tan", IntrinsicType::kTan},
IntrinsicData{"tanh", semantic::Intrinsic::kTanh}, IntrinsicData{"tanh", IntrinsicType::kTanh},
IntrinsicData{"textureDimensions", IntrinsicData{"textureDimensions", IntrinsicType::kTextureDimensions},
semantic::Intrinsic::kTextureDimensions}, IntrinsicData{"textureLoad", IntrinsicType::kTextureLoad},
IntrinsicData{"textureLoad", semantic::Intrinsic::kTextureLoad}, IntrinsicData{"textureNumLayers", IntrinsicType::kTextureNumLayers},
IntrinsicData{"textureNumLayers", IntrinsicData{"textureNumLevels", IntrinsicType::kTextureNumLevels},
semantic::Intrinsic::kTextureNumLayers}, IntrinsicData{"textureNumSamples", IntrinsicType::kTextureNumSamples},
IntrinsicData{"textureNumLevels", IntrinsicData{"textureSample", IntrinsicType::kTextureSample},
semantic::Intrinsic::kTextureNumLevels}, IntrinsicData{"textureSampleBias", IntrinsicType::kTextureSampleBias},
IntrinsicData{"textureNumSamples",
semantic::Intrinsic::kTextureNumSamples},
IntrinsicData{"textureSample", semantic::Intrinsic::kTextureSample},
IntrinsicData{"textureSampleBias",
semantic::Intrinsic::kTextureSampleBias},
IntrinsicData{"textureSampleCompare", IntrinsicData{"textureSampleCompare",
semantic::Intrinsic::kTextureSampleCompare}, IntrinsicType::kTextureSampleCompare},
IntrinsicData{"textureSampleGrad", IntrinsicData{"textureSampleGrad", IntrinsicType::kTextureSampleGrad},
semantic::Intrinsic::kTextureSampleGrad}, IntrinsicData{"textureSampleLevel", IntrinsicType::kTextureSampleLevel},
IntrinsicData{"textureSampleLevel", IntrinsicData{"trunc", IntrinsicType::kTrunc}));
semantic::Intrinsic::kTextureSampleLevel},
IntrinsicData{"trunc", semantic::Intrinsic::kTrunc}));
TEST_F(TypeDeterminerTest, MatchIntrinsicNoMatch) { TEST_F(TypeDeterminerTest, MatchIntrinsicNoMatch) {
EXPECT_EQ(TypeDeterminer::MatchIntrinsic("not_intrinsic"), EXPECT_EQ(TypeDeterminer::MatchIntrinsic("not_intrinsic"),
semantic::Intrinsic::kNone); IntrinsicType::kNone);
} }
using ImportData_DataPackingTest = TypeDeterminerTestWithParam<IntrinsicData>; using ImportData_DataPackingTest = TypeDeterminerTestWithParam<IntrinsicData>;
@ -1736,11 +1731,11 @@ INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest, TypeDeterminerTest,
ImportData_DataPackingTest, ImportData_DataPackingTest,
testing::Values( testing::Values(
IntrinsicData{"pack4x8snorm", semantic::Intrinsic::kPack4x8Snorm}, IntrinsicData{"pack4x8snorm", IntrinsicType::kPack4x8Snorm},
IntrinsicData{"pack4x8unorm", semantic::Intrinsic::kPack4x8Unorm}, IntrinsicData{"pack4x8unorm", IntrinsicType::kPack4x8Unorm},
IntrinsicData{"pack2x16snorm", semantic::Intrinsic::kPack2x16Snorm}, IntrinsicData{"pack2x16snorm", IntrinsicType::kPack2x16Snorm},
IntrinsicData{"pack2x16unorm", semantic::Intrinsic::kPack2x16Unorm}, IntrinsicData{"pack2x16unorm", IntrinsicType::kPack2x16Unorm},
IntrinsicData{"pack2x16float", semantic::Intrinsic::kPack2x16Float})); IntrinsicData{"pack2x16float", IntrinsicType::kPack2x16Float}));
using ImportData_SingleParamTest = TypeDeterminerTestWithParam<IntrinsicData>; using ImportData_SingleParamTest = TypeDeterminerTestWithParam<IntrinsicData>;
TEST_P(ImportData_SingleParamTest, Scalar) { TEST_P(ImportData_SingleParamTest, Scalar) {
@ -1786,29 +1781,28 @@ TEST_P(ImportData_SingleParamTest, Error_NoParams) {
INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest, TypeDeterminerTest,
ImportData_SingleParamTest, ImportData_SingleParamTest,
testing::Values(IntrinsicData{"acos", semantic::Intrinsic::kAcos}, testing::Values(IntrinsicData{"acos", IntrinsicType::kAcos},
IntrinsicData{"asin", semantic::Intrinsic::kAsin}, IntrinsicData{"asin", IntrinsicType::kAsin},
IntrinsicData{"atan", semantic::Intrinsic::kAtan}, IntrinsicData{"atan", IntrinsicType::kAtan},
IntrinsicData{"ceil", semantic::Intrinsic::kCeil}, IntrinsicData{"ceil", IntrinsicType::kCeil},
IntrinsicData{"cos", semantic::Intrinsic::kCos}, IntrinsicData{"cos", IntrinsicType::kCos},
IntrinsicData{"cosh", semantic::Intrinsic::kCosh}, IntrinsicData{"cosh", IntrinsicType::kCosh},
IntrinsicData{"exp", semantic::Intrinsic::kExp}, IntrinsicData{"exp", IntrinsicType::kExp},
IntrinsicData{"exp2", semantic::Intrinsic::kExp2}, IntrinsicData{"exp2", IntrinsicType::kExp2},
IntrinsicData{"floor", semantic::Intrinsic::kFloor}, IntrinsicData{"floor", IntrinsicType::kFloor},
IntrinsicData{"fract", semantic::Intrinsic::kFract}, IntrinsicData{"fract", IntrinsicType::kFract},
IntrinsicData{"inverseSqrt", IntrinsicData{"inverseSqrt", IntrinsicType::kInverseSqrt},
semantic::Intrinsic::kInverseSqrt}, IntrinsicData{"log", IntrinsicType::kLog},
IntrinsicData{"log", semantic::Intrinsic::kLog}, IntrinsicData{"log2", IntrinsicType::kLog2},
IntrinsicData{"log2", semantic::Intrinsic::kLog2}, IntrinsicData{"normalize", IntrinsicType::kNormalize},
IntrinsicData{"normalize", semantic::Intrinsic::kNormalize}, IntrinsicData{"round", IntrinsicType::kRound},
IntrinsicData{"round", semantic::Intrinsic::kRound}, IntrinsicData{"sign", IntrinsicType::kSign},
IntrinsicData{"sign", semantic::Intrinsic::kSign}, IntrinsicData{"sin", IntrinsicType::kSin},
IntrinsicData{"sin", semantic::Intrinsic::kSin}, IntrinsicData{"sinh", IntrinsicType::kSinh},
IntrinsicData{"sinh", semantic::Intrinsic::kSinh}, IntrinsicData{"sqrt", IntrinsicType::kSqrt},
IntrinsicData{"sqrt", semantic::Intrinsic::kSqrt}, IntrinsicData{"tan", IntrinsicType::kTan},
IntrinsicData{"tan", semantic::Intrinsic::kTan}, IntrinsicData{"tanh", IntrinsicType::kTanh},
IntrinsicData{"tanh", semantic::Intrinsic::kTanh}, IntrinsicData{"trunc", IntrinsicType::kTrunc}));
IntrinsicData{"trunc", semantic::Intrinsic::kTrunc}));
using ImportData_SingleParam_FloatOrInt_Test = using ImportData_SingleParam_FloatOrInt_Test =
TypeDeterminerTestWithParam<IntrinsicData>; TypeDeterminerTestWithParam<IntrinsicData>;
@ -1919,8 +1913,8 @@ TEST_P(ImportData_SingleParam_FloatOrInt_Test, Error_NoParams) {
INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest, INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest,
ImportData_SingleParam_FloatOrInt_Test, ImportData_SingleParam_FloatOrInt_Test,
testing::Values(IntrinsicData{ testing::Values(IntrinsicData{"abs",
"abs", semantic::Intrinsic::kAbs})); IntrinsicType::kAbs}));
TEST_F(TypeDeterminerTest, ImportData_Length_Scalar) { TEST_F(TypeDeterminerTest, ImportData_Length_Scalar) {
auto* ident = Expr("length"); auto* ident = Expr("length");
@ -1992,10 +1986,10 @@ TEST_P(ImportData_TwoParamTest, Error_NoParams) {
INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest, TypeDeterminerTest,
ImportData_TwoParamTest, ImportData_TwoParamTest,
testing::Values(IntrinsicData{"atan2", semantic::Intrinsic::kAtan2}, testing::Values(IntrinsicData{"atan2", IntrinsicType::kAtan2},
IntrinsicData{"pow", semantic::Intrinsic::kPow}, IntrinsicData{"pow", IntrinsicType::kPow},
IntrinsicData{"step", semantic::Intrinsic::kStep}, IntrinsicData{"step", IntrinsicType::kStep},
IntrinsicData{"reflect", semantic::Intrinsic::kReflect})); IntrinsicData{"reflect", IntrinsicType::kReflect}));
TEST_F(TypeDeterminerTest, ImportData_Distance_Scalar) { TEST_F(TypeDeterminerTest, ImportData_Distance_Scalar) {
auto* ident = Expr("distance"); auto* ident = Expr("distance");
@ -2093,11 +2087,10 @@ TEST_P(ImportData_ThreeParamTest, Error_NoParams) {
INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest, TypeDeterminerTest,
ImportData_ThreeParamTest, ImportData_ThreeParamTest,
testing::Values( testing::Values(IntrinsicData{"mix", IntrinsicType::kMix},
IntrinsicData{"mix", semantic::Intrinsic::kMix}, IntrinsicData{"smoothStep", IntrinsicType::kSmoothStep},
IntrinsicData{"smoothStep", semantic::Intrinsic::kSmoothStep}, IntrinsicData{"fma", IntrinsicType::kFma},
IntrinsicData{"fma", semantic::Intrinsic::kFma}, IntrinsicData{"faceForward", IntrinsicType::kFaceForward}));
IntrinsicData{"faceForward", semantic::Intrinsic::kFaceForward}));
using ImportData_ThreeParam_FloatOrInt_Test = using ImportData_ThreeParam_FloatOrInt_Test =
TypeDeterminerTestWithParam<IntrinsicData>; TypeDeterminerTestWithParam<IntrinsicData>;
@ -2200,8 +2193,8 @@ TEST_P(ImportData_ThreeParam_FloatOrInt_Test, Error_NoParams) {
INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest, INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest,
ImportData_ThreeParam_FloatOrInt_Test, ImportData_ThreeParam_FloatOrInt_Test,
testing::Values(IntrinsicData{ testing::Values(IntrinsicData{"clamp",
"clamp", semantic::Intrinsic::kClamp})); IntrinsicType::kClamp}));
using ImportData_Int_SingleParamTest = using ImportData_Int_SingleParamTest =
TypeDeterminerTestWithParam<IntrinsicData>; TypeDeterminerTestWithParam<IntrinsicData>;
@ -2248,9 +2241,8 @@ TEST_P(ImportData_Int_SingleParamTest, Error_NoParams) {
INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest, TypeDeterminerTest,
ImportData_Int_SingleParamTest, ImportData_Int_SingleParamTest,
testing::Values( testing::Values(IntrinsicData{"countOneBits", IntrinsicType::kCountOneBits},
IntrinsicData{"countOneBits", semantic::Intrinsic::kCountOneBits}, IntrinsicData{"reverseBits", IntrinsicType::kReverseBits}));
IntrinsicData{"reverseBits", semantic::Intrinsic::kReverseBits}));
using ImportData_FloatOrInt_TwoParamTest = using ImportData_FloatOrInt_TwoParamTest =
TypeDeterminerTestWithParam<IntrinsicData>; TypeDeterminerTestWithParam<IntrinsicData>;
@ -2351,8 +2343,8 @@ TEST_P(ImportData_FloatOrInt_TwoParamTest, Error_NoParams) {
INSTANTIATE_TEST_SUITE_P( INSTANTIATE_TEST_SUITE_P(
TypeDeterminerTest, TypeDeterminerTest,
ImportData_FloatOrInt_TwoParamTest, ImportData_FloatOrInt_TwoParamTest,
testing::Values(IntrinsicData{"min", semantic::Intrinsic::kMin}, testing::Values(IntrinsicData{"min", IntrinsicType::kMin},
IntrinsicData{"max", semantic::Intrinsic::kMax})); IntrinsicData{"max", IntrinsicType::kMax}));
TEST_F(TypeDeterminerTest, ImportData_GLSL_Determinant) { TEST_F(TypeDeterminerTest, ImportData_GLSL_Determinant) {
Global("var", ast::StorageClass::kFunction, ty.mat3x3<f32>()); Global("var", ast::StorageClass::kFunction, ty.mat3x3<f32>());
@ -2385,8 +2377,7 @@ TEST_P(ImportData_Matrix_OneParam_Test, NoParams) {
INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest, INSTANTIATE_TEST_SUITE_P(TypeDeterminerTest,
ImportData_Matrix_OneParam_Test, ImportData_Matrix_OneParam_Test,
testing::Values(IntrinsicData{ testing::Values(IntrinsicData{
"determinant", "determinant", IntrinsicType::kDeterminant}));
semantic::Intrinsic::kDeterminant}));
TEST_F(TypeDeterminerTest, Function_EntryPoints_StageDecoration) { TEST_F(TypeDeterminerTest, Function_EntryPoints_StageDecoration) {
// fn b() {} // fn b() {}

197
src/validator/validator_impl.cc
View File

@ -49,6 +49,8 @@ namespace tint {
namespace { namespace {
using IntrinsicType = semantic::IntrinsicType;
enum class IntrinsicDataType { enum class IntrinsicDataType {
kMixed, kMixed,
kFloatOrIntScalarOrVector, kFloatOrIntScalarOrVector,
@ -63,7 +65,7 @@ enum class IntrinsicDataType {
}; };
struct IntrinsicData { struct IntrinsicData {
semantic::Intrinsic intrinsic; IntrinsicType intrinsic;
uint32_t param_count; uint32_t param_count;
IntrinsicDataType data_type; IntrinsicDataType data_type;
uint32_t vector_size; uint32_t vector_size;
@ -73,122 +75,97 @@ struct IntrinsicData {
// Note, this isn't all the intrinsics. Some are handled specially before // Note, this isn't all the intrinsics. Some are handled specially before
// we get to the generic code. See the ValidateCallExpr code below. // we get to the generic code. See the ValidateCallExpr code below.
constexpr const IntrinsicData kIntrinsicData[] = { constexpr const IntrinsicData kIntrinsicData[] = {
{semantic::Intrinsic::kAbs, 1, IntrinsicDataType::kFloatOrIntScalarOrVector, {IntrinsicType::kAbs, 1, IntrinsicDataType::kFloatOrIntScalarOrVector, 0,
0, true},
{semantic::Intrinsic::kAcos, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true}, true},
{semantic::Intrinsic::kAll, 1, IntrinsicDataType::kBoolVector, 0, false}, {IntrinsicType::kAcos, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
{semantic::Intrinsic::kAny, 1, IntrinsicDataType::kBoolVector, 0, false}, {IntrinsicType::kAll, 1, IntrinsicDataType::kBoolVector, 0, false},
{semantic::Intrinsic::kArrayLength, 1, IntrinsicDataType::kMixed, 0, false}, {IntrinsicType::kAny, 1, IntrinsicDataType::kBoolVector, 0, false},
{semantic::Intrinsic::kAsin, 1, IntrinsicDataType::kFloatScalarOrVector, 0, {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}, true},
{semantic::Intrinsic::kAtan, 1, IntrinsicDataType::kFloatScalarOrVector, 0, {IntrinsicType::kCeil, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
{IntrinsicType::kClamp, 3, IntrinsicDataType::kFloatOrIntScalarOrVector, 0,
true}, 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}, true},
{semantic::Intrinsic::kCeil, 1, IntrinsicDataType::kFloatScalarOrVector, 0, {IntrinsicType::kCross, 2, IntrinsicDataType::kFloatVector, 3, true},
true}, {IntrinsicType::kDeterminant, 1, IntrinsicDataType::kMatrix, 0, false},
{semantic::Intrinsic::kClamp, 3, {IntrinsicType::kDistance, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
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,
false}, false},
{semantic::Intrinsic::kDistance, 2, IntrinsicDataType::kFloatScalarOrVector, {IntrinsicType::kDot, 2, IntrinsicDataType::kFloatVector, 0, false},
0, false}, {IntrinsicType::kDpdx, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
{semantic::Intrinsic::kDot, 2, IntrinsicDataType::kFloatVector, 0, false}, {IntrinsicType::kDpdxCoarse, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
{semantic::Intrinsic::kDpdx, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true}, true},
{semantic::Intrinsic::kDpdxCoarse, 1, {IntrinsicType::kDpdxFine, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
IntrinsicDataType::kFloatScalarOrVector, 0, true}, true},
{semantic::Intrinsic::kDpdxFine, 1, IntrinsicDataType::kFloatScalarOrVector, {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}, 0, true},
{semantic::Intrinsic::kDpdy, 1, IntrinsicDataType::kFloatScalarOrVector, 0, {IntrinsicType::kFwidthFine, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true}, true},
{semantic::Intrinsic::kDpdyCoarse, 1, {IntrinsicType::kInverseSqrt, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
IntrinsicDataType::kFloatScalarOrVector, 0, true},
{semantic::Intrinsic::kDpdyFine, 1, IntrinsicDataType::kFloatScalarOrVector,
0, true},
{semantic::Intrinsic::kExp, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true}, true},
{semantic::Intrinsic::kExp2, 1, IntrinsicDataType::kFloatScalarOrVector, 0, {IntrinsicType::kLdexp, 2, IntrinsicDataType::kFloatScalarOrVector, 0,
true}, true},
{semantic::Intrinsic::kFaceForward, 3, {IntrinsicType::kLength, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
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,
false}, 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}, false},
{semantic::Intrinsic::kPack2x16Snorm, 1, IntrinsicDataType::kFloatVector, 2, {IntrinsicType::kPack4x8Unorm, 1, IntrinsicDataType::kFloatVector, 4,
false}, false},
{semantic::Intrinsic::kPack2x16Unorm, 1, IntrinsicDataType::kFloatVector, 2, {IntrinsicType::kPack2x16Snorm, 1, IntrinsicDataType::kFloatVector, 2,
false}, false},
{semantic::Intrinsic::kPack2x16Float, 1, IntrinsicDataType::kFloatVector, 2, {IntrinsicType::kPack2x16Unorm, 1, IntrinsicDataType::kFloatVector, 2,
false}, 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}, true},
{semantic::Intrinsic::kReflect, 2, IntrinsicDataType::kFloatScalarOrVector, {IntrinsicType::kReverseBits, 1, IntrinsicDataType::kIntScalarOrVector, 0,
0, true},
{semantic::Intrinsic::kReverseBits, 1,
IntrinsicDataType::kIntScalarOrVector, 0, true},
{semantic::Intrinsic::kRound, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true}, true},
{semantic::Intrinsic::kSelect, 3, IntrinsicDataType::kMixed, 0, false}, {IntrinsicType::kRound, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
{semantic::Intrinsic::kSign, 1, IntrinsicDataType::kFloatScalarOrVector, 0,
true}, 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}, true},
{semantic::Intrinsic::kSinh, 1, IntrinsicDataType::kFloatScalarOrVector, 0, {IntrinsicType::kSqrt, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
true}, {IntrinsicType::kStep, 2, IntrinsicDataType::kFloatScalarOrVector, 0, true},
{semantic::Intrinsic::kSmoothStep, 3, {IntrinsicType::kTan, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
IntrinsicDataType::kFloatScalarOrVector, 0, true}, {IntrinsicType::kTanh, 1, IntrinsicDataType::kFloatScalarOrVector, 0, true},
{semantic::Intrinsic::kSqrt, 1, IntrinsicDataType::kFloatScalarOrVector, 0, {IntrinsicType::kTrunc, 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,
true}, true},
}; };
@ -731,7 +708,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
} }
} else { } else {
// Special cases. // Special cases.
if (data->intrinsic == semantic::Intrinsic::kFrexp) { if (data->intrinsic == IntrinsicType::kFrexp) {
auto* p0 = expr->params()[0]; auto* p0 = expr->params()[0];
auto* p1 = expr->params()[1]; auto* p1 = expr->params()[1];
auto* t0 = program_->TypeOf(p0)->UnwrapPtrIfNeeded(); auto* t0 = program_->TypeOf(p0)->UnwrapPtrIfNeeded();
@ -773,7 +750,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
} }
} }
if (data->intrinsic == semantic::Intrinsic::kSelect) { if (data->intrinsic == IntrinsicType::kSelect) {
auto* type = program_->TypeOf(expr); auto* type = program_->TypeOf(expr);
auto* t0 = program_->TypeOf(expr->params()[0])->UnwrapPtrIfNeeded(); auto* t0 = program_->TypeOf(expr->params()[0])->UnwrapPtrIfNeeded();
auto* t1 = program_->TypeOf(expr->params()[1])->UnwrapPtrIfNeeded(); auto* t1 = program_->TypeOf(expr->params()[1])->UnwrapPtrIfNeeded();
@ -819,7 +796,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
} }
} }
if (data->intrinsic == semantic::Intrinsic::kArrayLength) { if (data->intrinsic == IntrinsicType::kArrayLength) {
if (!program_->TypeOf(expr)->UnwrapPtrIfNeeded()->Is<type::U32>()) { if (!program_->TypeOf(expr)->UnwrapPtrIfNeeded()->Is<type::U32>()) {
add_error( add_error(
expr->source(), expr->source(),
@ -840,15 +817,15 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
} }
// Result types don't match parameter types. // Result types don't match parameter types.
if (data->intrinsic == semantic::Intrinsic::kAll || if (data->intrinsic == IntrinsicType::kAll ||
data->intrinsic == semantic::Intrinsic::kAny) { data->intrinsic == IntrinsicType::kAny) {
if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin, if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
IntrinsicDataType::kBoolScalar, 0, this)) { IntrinsicDataType::kBoolScalar, 0, this)) {
return false; return false;
} }
} }
if (data->intrinsic == semantic::Intrinsic::kDot) { if (data->intrinsic == IntrinsicType::kDot) {
if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin, if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
IntrinsicDataType::kFloatScalar, 0, this)) { IntrinsicDataType::kFloatScalar, 0, this)) {
return false; return false;
@ -864,9 +841,9 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
} }
} }
if (data->intrinsic == semantic::Intrinsic::kLength || if (data->intrinsic == IntrinsicType::kLength ||
data->intrinsic == semantic::Intrinsic::kDistance || data->intrinsic == IntrinsicType::kDistance ||
data->intrinsic == semantic::Intrinsic::kDeterminant) { data->intrinsic == IntrinsicType::kDeterminant) {
if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin, if (!IsValidType(program_->TypeOf(expr), expr->source(), builtin,
IntrinsicDataType::kFloatScalar, 0, this)) { IntrinsicDataType::kFloatScalar, 0, this)) {
return false; return false;
@ -874,7 +851,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
} }
// Must be a square matrix. // Must be a square matrix.
if (data->intrinsic == semantic::Intrinsic::kDeterminant) { if (data->intrinsic == IntrinsicType::kDeterminant) {
const auto* matrix = const auto* matrix =
program_->TypeOf(expr->params()[0])->As<type::Matrix>(); program_->TypeOf(expr->params()[0])->As<type::Matrix>();
if (matrix->rows() != matrix->columns()) { if (matrix->rows() != matrix->columns()) {
@ -887,8 +864,8 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
} }
// Last parameter must be a pointer. // Last parameter must be a pointer.
if (data->intrinsic == semantic::Intrinsic::kFrexp || if (data->intrinsic == IntrinsicType::kFrexp ||
data->intrinsic == semantic::Intrinsic::kModf) { data->intrinsic == IntrinsicType::kModf) {
auto* last_param = expr->params()[data->param_count - 1]; auto* last_param = expr->params()[data->param_count - 1];
if (!program_->TypeOf(last_param)->Is<type::Pointer>()) { if (!program_->TypeOf(last_param)->Is<type::Pointer>()) {
add_error(last_param->source(), "incorrect type for " + builtin + add_error(last_param->source(), "incorrect type for " + builtin +

166
src/writer/hlsl/generator_impl.cc
View File

@ -548,10 +548,10 @@ bool GeneratorImpl::EmitCall(std::ostream& pre,
} }
if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) { if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) {
const auto& params = expr->params(); 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"; error_ = "select not supported in HLSL backend yet";
return false; 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"; error_ = "is_normal not supported in HLSL backend yet";
return false; return false;
} else if (semantic::intrinsic::IsDataPackingIntrinsic(sem->intrinsic())) { } else if (semantic::intrinsic::IsDataPackingIntrinsic(sem->intrinsic())) {
@ -650,21 +650,21 @@ bool GeneratorImpl::EmitDataPackingCall(std::ostream& pre,
uint32_t dims = 2; uint32_t dims = 2;
bool is_signed = false; bool is_signed = false;
uint32_t scale = 65535; uint32_t scale = 65535;
if (ident->intrinsic() == semantic::Intrinsic::kPack4x8Snorm || if (ident->intrinsic() == semantic::IntrinsicType::kPack4x8Snorm ||
ident->intrinsic() == semantic::Intrinsic::kPack4x8Unorm) { ident->intrinsic() == semantic::IntrinsicType::kPack4x8Unorm) {
dims = 4; dims = 4;
scale = 255; scale = 255;
} }
if (ident->intrinsic() == semantic::Intrinsic::kPack4x8Snorm || if (ident->intrinsic() == semantic::IntrinsicType::kPack4x8Snorm ||
ident->intrinsic() == semantic::Intrinsic::kPack2x16Snorm) { ident->intrinsic() == semantic::IntrinsicType::kPack2x16Snorm) {
is_signed = true; is_signed = true;
scale = (scale - 1) / 2; scale = (scale - 1) / 2;
} }
switch (ident->intrinsic()) { switch (ident->intrinsic()) {
case semantic::Intrinsic::kPack4x8Snorm: case semantic::IntrinsicType::kPack4x8Snorm:
case semantic::Intrinsic::kPack4x8Unorm: case semantic::IntrinsicType::kPack4x8Unorm:
case semantic::Intrinsic::kPack2x16Snorm: case semantic::IntrinsicType::kPack2x16Snorm:
case semantic::Intrinsic::kPack2x16Unorm: case semantic::IntrinsicType::kPack2x16Unorm:
pre << (is_signed ? "" : "u") << "int" << dims << " " << tmp_name << " = " pre << (is_signed ? "" : "u") << "int" << dims << " " << tmp_name << " = "
<< (is_signed ? "" : "u") << "int" << dims << "(round(clamp(" << (is_signed ? "" : "u") << "int" << dims << "(round(clamp("
<< expr_out.str() << ", " << (is_signed ? "-1.0" : "0.0") << expr_out.str() << ", " << (is_signed ? "-1.0" : "0.0")
@ -683,7 +683,7 @@ bool GeneratorImpl::EmitDataPackingCall(std::ostream& pre,
} }
out << ")"; out << ")";
break; break;
case semantic::Intrinsic::kPack2x16Float: case semantic::IntrinsicType::kPack2x16Float:
pre << "uint2 " << tmp_name << " = f32tof16(" << expr_out.str() << ");\n"; pre << "uint2 " << tmp_name << " = f32tof16(" << expr_out.str() << ");\n";
out << "(" << tmp_name << ".x | " << tmp_name << ".y << 16)"; out << "(" << tmp_name << ".x | " << tmp_name << ".y << 16)";
break; break;
@ -709,17 +709,17 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
auto* texture_type = TypeOf(texture)->UnwrapAll()->As<type::Texture>(); auto* texture_type = TypeOf(texture)->UnwrapAll()->As<type::Texture>();
switch (sem->intrinsic()) { switch (sem->intrinsic()) {
case semantic::Intrinsic::kTextureDimensions: case semantic::IntrinsicType::kTextureDimensions:
case semantic::Intrinsic::kTextureNumLayers: case semantic::IntrinsicType::kTextureNumLayers:
case semantic::Intrinsic::kTextureNumLevels: case semantic::IntrinsicType::kTextureNumLevels:
case semantic::Intrinsic::kTextureNumSamples: { case semantic::IntrinsicType::kTextureNumSamples: {
// All of these intrinsics use the GetDimensions() method on the texture // All of these intrinsics use the GetDimensions() method on the texture
int num_dimensions = 0; int num_dimensions = 0;
const char* swizzle = ""; const char* swizzle = "";
bool add_mip_level_in = false; bool add_mip_level_in = false;
switch (sem->intrinsic()) { switch (sem->intrinsic()) {
case semantic::Intrinsic::kTextureDimensions: case semantic::IntrinsicType::kTextureDimensions:
switch (texture_type->dim()) { switch (texture_type->dim()) {
case type::TextureDimension::kNone: case type::TextureDimension::kNone:
error_ = "texture dimension is kNone"; error_ = "texture dimension is kNone";
@ -755,7 +755,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
break; break;
} }
break; break;
case semantic::Intrinsic::kTextureNumLayers: case semantic::IntrinsicType::kTextureNumLayers:
switch (texture_type->dim()) { switch (texture_type->dim()) {
default: default:
error_ = "texture dimension is not arrayed"; error_ = "texture dimension is not arrayed";
@ -771,7 +771,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
break; break;
} }
break; break;
case semantic::Intrinsic::kTextureNumLevels: case semantic::IntrinsicType::kTextureNumLevels:
add_mip_level_in = true; add_mip_level_in = true;
switch (texture_type->dim()) { switch (texture_type->dim()) {
default: default:
@ -790,7 +790,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
break; break;
} }
break; break;
case semantic::Intrinsic::kTextureNumSamples: case semantic::IntrinsicType::kTextureNumSamples:
switch (texture_type->dim()) { switch (texture_type->dim()) {
default: default:
error_ = "texture dimension does not support multisampling"; error_ = "texture dimension does not support multisampling";
@ -860,28 +860,28 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
bool pack_mip_in_coords = false; bool pack_mip_in_coords = false;
switch (sem->intrinsic()) { switch (sem->intrinsic()) {
case semantic::Intrinsic::kTextureSample: case semantic::IntrinsicType::kTextureSample:
out << ".Sample("; out << ".Sample(";
break; break;
case semantic::Intrinsic::kTextureSampleBias: case semantic::IntrinsicType::kTextureSampleBias:
out << ".SampleBias("; out << ".SampleBias(";
break; break;
case semantic::Intrinsic::kTextureSampleLevel: case semantic::IntrinsicType::kTextureSampleLevel:
out << ".SampleLevel("; out << ".SampleLevel(";
break; break;
case semantic::Intrinsic::kTextureSampleGrad: case semantic::IntrinsicType::kTextureSampleGrad:
out << ".SampleGrad("; out << ".SampleGrad(";
break; break;
case semantic::Intrinsic::kTextureSampleCompare: case semantic::IntrinsicType::kTextureSampleCompare:
out << ".SampleCmp("; out << ".SampleCmp(";
break; break;
case semantic::Intrinsic::kTextureLoad: case semantic::IntrinsicType::kTextureLoad:
out << ".Load("; out << ".Load(";
if (!texture_type->Is<type::StorageTexture>()) { if (!texture_type->Is<type::StorageTexture>()) {
pack_mip_in_coords = true; pack_mip_in_coords = true;
} }
break; break;
case semantic::Intrinsic::kTextureStore: case semantic::IntrinsicType::kTextureStore:
out << "["; out << "[";
break; break;
default: default:
@ -937,7 +937,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre,
} }
} }
if (sem->intrinsic() == semantic::Intrinsic::kTextureStore) { if (sem->intrinsic() == semantic::IntrinsicType::kTextureStore) {
out << "] = "; out << "] = ";
if (!EmitExpression(pre, out, params[pidx.value])) if (!EmitExpression(pre, out, params[pidx.value]))
return false; return false;
@ -952,94 +952,94 @@ std::string GeneratorImpl::generate_builtin_name(
const semantic::IntrinsicCall* call) { const semantic::IntrinsicCall* call) {
std::string out; std::string out;
switch (call->intrinsic()) { switch (call->intrinsic()) {
case semantic::Intrinsic::kAcos: case semantic::IntrinsicType::kAcos:
case semantic::Intrinsic::kAny: case semantic::IntrinsicType::kAny:
case semantic::Intrinsic::kAll: case semantic::IntrinsicType::kAll:
case semantic::Intrinsic::kAsin: case semantic::IntrinsicType::kAsin:
case semantic::Intrinsic::kAtan: case semantic::IntrinsicType::kAtan:
case semantic::Intrinsic::kAtan2: case semantic::IntrinsicType::kAtan2:
case semantic::Intrinsic::kCeil: case semantic::IntrinsicType::kCeil:
case semantic::Intrinsic::kCos: case semantic::IntrinsicType::kCos:
case semantic::Intrinsic::kCosh: case semantic::IntrinsicType::kCosh:
case semantic::Intrinsic::kCross: case semantic::IntrinsicType::kCross:
case semantic::Intrinsic::kDeterminant: case semantic::IntrinsicType::kDeterminant:
case semantic::Intrinsic::kDistance: case semantic::IntrinsicType::kDistance:
case semantic::Intrinsic::kDot: case semantic::IntrinsicType::kDot:
case semantic::Intrinsic::kExp: case semantic::IntrinsicType::kExp:
case semantic::Intrinsic::kExp2: case semantic::IntrinsicType::kExp2:
case semantic::Intrinsic::kFloor: case semantic::IntrinsicType::kFloor:
case semantic::Intrinsic::kFma: case semantic::IntrinsicType::kFma:
case semantic::Intrinsic::kLdexp: case semantic::IntrinsicType::kLdexp:
case semantic::Intrinsic::kLength: case semantic::IntrinsicType::kLength:
case semantic::Intrinsic::kLog: case semantic::IntrinsicType::kLog:
case semantic::Intrinsic::kLog2: case semantic::IntrinsicType::kLog2:
case semantic::Intrinsic::kNormalize: case semantic::IntrinsicType::kNormalize:
case semantic::Intrinsic::kPow: case semantic::IntrinsicType::kPow:
case semantic::Intrinsic::kReflect: case semantic::IntrinsicType::kReflect:
case semantic::Intrinsic::kRound: case semantic::IntrinsicType::kRound:
case semantic::Intrinsic::kSin: case semantic::IntrinsicType::kSin:
case semantic::Intrinsic::kSinh: case semantic::IntrinsicType::kSinh:
case semantic::Intrinsic::kSqrt: case semantic::IntrinsicType::kSqrt:
case semantic::Intrinsic::kStep: case semantic::IntrinsicType::kStep:
case semantic::Intrinsic::kTan: case semantic::IntrinsicType::kTan:
case semantic::Intrinsic::kTanh: case semantic::IntrinsicType::kTanh:
case semantic::Intrinsic::kTrunc: case semantic::IntrinsicType::kTrunc:
case semantic::Intrinsic::kMix: case semantic::IntrinsicType::kMix:
case semantic::Intrinsic::kSign: case semantic::IntrinsicType::kSign:
case semantic::Intrinsic::kAbs: case semantic::IntrinsicType::kAbs:
case semantic::Intrinsic::kMax: case semantic::IntrinsicType::kMax:
case semantic::Intrinsic::kMin: case semantic::IntrinsicType::kMin:
case semantic::Intrinsic::kClamp: case semantic::IntrinsicType::kClamp:
out = semantic::intrinsic::str(call->intrinsic()); out = semantic::intrinsic::str(call->intrinsic());
break; break;
case semantic::Intrinsic::kCountOneBits: case semantic::IntrinsicType::kCountOneBits:
out = "countbits"; out = "countbits";
break; break;
case semantic::Intrinsic::kDpdx: case semantic::IntrinsicType::kDpdx:
out = "ddx"; out = "ddx";
break; break;
case semantic::Intrinsic::kDpdxCoarse: case semantic::IntrinsicType::kDpdxCoarse:
out = "ddx_coarse"; out = "ddx_coarse";
break; break;
case semantic::Intrinsic::kDpdxFine: case semantic::IntrinsicType::kDpdxFine:
out = "ddx_fine"; out = "ddx_fine";
break; break;
case semantic::Intrinsic::kDpdy: case semantic::IntrinsicType::kDpdy:
out = "ddy"; out = "ddy";
break; break;
case semantic::Intrinsic::kDpdyCoarse: case semantic::IntrinsicType::kDpdyCoarse:
out = "ddy_coarse"; out = "ddy_coarse";
break; break;
case semantic::Intrinsic::kDpdyFine: case semantic::IntrinsicType::kDpdyFine:
out = "ddy_fine"; out = "ddy_fine";
break; break;
case semantic::Intrinsic::kFaceForward: case semantic::IntrinsicType::kFaceForward:
out = "faceforward"; out = "faceforward";
break; break;
case semantic::Intrinsic::kFract: case semantic::IntrinsicType::kFract:
out = "frac"; out = "frac";
break; break;
case semantic::Intrinsic::kFwidth: case semantic::IntrinsicType::kFwidth:
case semantic::Intrinsic::kFwidthCoarse: case semantic::IntrinsicType::kFwidthCoarse:
case semantic::Intrinsic::kFwidthFine: case semantic::IntrinsicType::kFwidthFine:
out = "fwidth"; out = "fwidth";
break; break;
case semantic::Intrinsic::kInverseSqrt: case semantic::IntrinsicType::kInverseSqrt:
out = "rsqrt"; out = "rsqrt";
break; break;
case semantic::Intrinsic::kIsFinite: case semantic::IntrinsicType::kIsFinite:
out = "isfinite"; out = "isfinite";
break; break;
case semantic::Intrinsic::kIsInf: case semantic::IntrinsicType::kIsInf:
out = "isinf"; out = "isinf";
break; break;
case semantic::Intrinsic::kIsNan: case semantic::IntrinsicType::kIsNan:
out = "isnan"; out = "isnan";
break; break;
case semantic::Intrinsic::kReverseBits: case semantic::IntrinsicType::kReverseBits:
out = "reversebits"; out = "reversebits";
break; break;
case semantic::Intrinsic::kSmoothStep: case semantic::IntrinsicType::kSmoothStep:
out = "smoothstep"; out = "smoothstep";
break; break;
default: default:

249
src/writer/hlsl/generator_impl_intrinsic_test.cc
View File

@ -29,6 +29,8 @@ namespace writer {
namespace hlsl { namespace hlsl {
namespace { namespace {
using IntrinsicType = semantic::IntrinsicType;
using ::testing::HasSubstr; using ::testing::HasSubstr;
using HlslGeneratorImplTest_Intrinsic = TestHelper; using HlslGeneratorImplTest_Intrinsic = TestHelper;
@ -40,7 +42,7 @@ enum class ParamType {
}; };
struct IntrinsicData { struct IntrinsicData {
semantic::Intrinsic intrinsic; IntrinsicType intrinsic;
ParamType type; ParamType type;
const char* hlsl_name; const char* hlsl_name;
}; };
@ -61,92 +63,92 @@ inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
return out; return out;
} }
ast::CallExpression* GenerateCall(semantic::Intrinsic intrinsic, ast::CallExpression* GenerateCall(IntrinsicType intrinsic,
ParamType type, ParamType type,
ProgramBuilder* builder) { ProgramBuilder* builder) {
std::string name; std::string name;
std::ostringstream str(name); std::ostringstream str(name);
str << intrinsic; str << intrinsic;
switch (intrinsic) { switch (intrinsic) {
case semantic::Intrinsic::kAcos: case IntrinsicType::kAcos:
case semantic::Intrinsic::kAsin: case IntrinsicType::kAsin:
case semantic::Intrinsic::kAtan: case IntrinsicType::kAtan:
case semantic::Intrinsic::kCeil: case IntrinsicType::kCeil:
case semantic::Intrinsic::kCos: case IntrinsicType::kCos:
case semantic::Intrinsic::kCosh: case IntrinsicType::kCosh:
case semantic::Intrinsic::kDpdx: case IntrinsicType::kDpdx:
case semantic::Intrinsic::kDpdxCoarse: case IntrinsicType::kDpdxCoarse:
case semantic::Intrinsic::kDpdxFine: case IntrinsicType::kDpdxFine:
case semantic::Intrinsic::kDpdy: case IntrinsicType::kDpdy:
case semantic::Intrinsic::kDpdyCoarse: case IntrinsicType::kDpdyCoarse:
case semantic::Intrinsic::kDpdyFine: case IntrinsicType::kDpdyFine:
case semantic::Intrinsic::kExp: case IntrinsicType::kExp:
case semantic::Intrinsic::kExp2: case IntrinsicType::kExp2:
case semantic::Intrinsic::kFloor: case IntrinsicType::kFloor:
case semantic::Intrinsic::kFract: case IntrinsicType::kFract:
case semantic::Intrinsic::kFwidth: case IntrinsicType::kFwidth:
case semantic::Intrinsic::kFwidthCoarse: case IntrinsicType::kFwidthCoarse:
case semantic::Intrinsic::kFwidthFine: case IntrinsicType::kFwidthFine:
case semantic::Intrinsic::kInverseSqrt: case IntrinsicType::kInverseSqrt:
case semantic::Intrinsic::kIsFinite: case IntrinsicType::kIsFinite:
case semantic::Intrinsic::kIsInf: case IntrinsicType::kIsInf:
case semantic::Intrinsic::kIsNan: case IntrinsicType::kIsNan:
case semantic::Intrinsic::kIsNormal: case IntrinsicType::kIsNormal:
case semantic::Intrinsic::kLdexp: case IntrinsicType::kLdexp:
case semantic::Intrinsic::kLength: case IntrinsicType::kLength:
case semantic::Intrinsic::kLog: case IntrinsicType::kLog:
case semantic::Intrinsic::kLog2: case IntrinsicType::kLog2:
case semantic::Intrinsic::kNormalize: case IntrinsicType::kNormalize:
case semantic::Intrinsic::kReflect: case IntrinsicType::kReflect:
case semantic::Intrinsic::kRound: case IntrinsicType::kRound:
case semantic::Intrinsic::kSin: case IntrinsicType::kSin:
case semantic::Intrinsic::kSinh: case IntrinsicType::kSinh:
case semantic::Intrinsic::kSqrt: case IntrinsicType::kSqrt:
case semantic::Intrinsic::kTan: case IntrinsicType::kTan:
case semantic::Intrinsic::kTanh: case IntrinsicType::kTanh:
case semantic::Intrinsic::kTrunc: case IntrinsicType::kTrunc:
case semantic::Intrinsic::kSign: case IntrinsicType::kSign:
return builder->Call(str.str(), "f1"); return builder->Call(str.str(), "f1");
case semantic::Intrinsic::kAtan2: case IntrinsicType::kAtan2:
case semantic::Intrinsic::kCross: case IntrinsicType::kCross:
case semantic::Intrinsic::kDot: case IntrinsicType::kDot:
case semantic::Intrinsic::kDistance: case IntrinsicType::kDistance:
case semantic::Intrinsic::kPow: case IntrinsicType::kPow:
case semantic::Intrinsic::kStep: case IntrinsicType::kStep:
return builder->Call(str.str(), "f1", "f2"); return builder->Call(str.str(), "f1", "f2");
case semantic::Intrinsic::kFma: case IntrinsicType::kFma:
case semantic::Intrinsic::kMix: case IntrinsicType::kMix:
case semantic::Intrinsic::kFaceForward: case IntrinsicType::kFaceForward:
case semantic::Intrinsic::kSmoothStep: case IntrinsicType::kSmoothStep:
return builder->Call(str.str(), "f1", "f2", "f3"); return builder->Call(str.str(), "f1", "f2", "f3");
case semantic::Intrinsic::kAll: case IntrinsicType::kAll:
case semantic::Intrinsic::kAny: case IntrinsicType::kAny:
return builder->Call(str.str(), "b1"); return builder->Call(str.str(), "b1");
case semantic::Intrinsic::kAbs: case IntrinsicType::kAbs:
if (type == ParamType::kF32) { if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1"); return builder->Call(str.str(), "f1");
} else { } else {
return builder->Call(str.str(), "u1"); return builder->Call(str.str(), "u1");
} }
case semantic::Intrinsic::kCountOneBits: case IntrinsicType::kCountOneBits:
case semantic::Intrinsic::kReverseBits: case IntrinsicType::kReverseBits:
return builder->Call(str.str(), "u1"); return builder->Call(str.str(), "u1");
case semantic::Intrinsic::kMax: case IntrinsicType::kMax:
case semantic::Intrinsic::kMin: case IntrinsicType::kMin:
if (type == ParamType::kF32) { if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1", "f2"); return builder->Call(str.str(), "f1", "f2");
} else { } else {
return builder->Call(str.str(), "u1", "u2"); return builder->Call(str.str(), "u1", "u2");
} }
case semantic::Intrinsic::kClamp: case IntrinsicType::kClamp:
if (type == ParamType::kF32) { if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1", "f2", "f3"); return builder->Call(str.str(), "f1", "f2", "f3");
} else { } else {
return builder->Call(str.str(), "u1", "u2", "u3"); return builder->Call(str.str(), "u1", "u2", "u3");
} }
case semantic::Intrinsic::kSelect: case IntrinsicType::kSelect:
return builder->Call(str.str(), "f1", "f2", "b1"); return builder->Call(str.str(), "f1", "f2", "b1");
case semantic::Intrinsic::kDeterminant: case IntrinsicType::kDeterminant:
return builder->Call(str.str(), "m1"); return builder->Call(str.str(), "m1");
default: default:
break; break;
@ -183,81 +185,72 @@ INSTANTIATE_TEST_SUITE_P(
HlslGeneratorImplTest_Intrinsic, HlslGeneratorImplTest_Intrinsic,
HlslIntrinsicTest, HlslIntrinsicTest,
testing::Values( testing::Values(
IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kF32, "abs"}, IntrinsicData{IntrinsicType::kAbs, ParamType::kF32, "abs"},
IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kU32, "abs"}, IntrinsicData{IntrinsicType::kAbs, ParamType::kU32, "abs"},
IntrinsicData{semantic::Intrinsic::kAcos, ParamType::kF32, "acos"}, IntrinsicData{IntrinsicType::kAcos, ParamType::kF32, "acos"},
IntrinsicData{semantic::Intrinsic::kAll, ParamType::kBool, "all"}, IntrinsicData{IntrinsicType::kAll, ParamType::kBool, "all"},
IntrinsicData{semantic::Intrinsic::kAny, ParamType::kBool, "any"}, IntrinsicData{IntrinsicType::kAny, ParamType::kBool, "any"},
IntrinsicData{semantic::Intrinsic::kAsin, ParamType::kF32, "asin"}, IntrinsicData{IntrinsicType::kAsin, ParamType::kF32, "asin"},
IntrinsicData{semantic::Intrinsic::kAtan, ParamType::kF32, "atan"}, IntrinsicData{IntrinsicType::kAtan, ParamType::kF32, "atan"},
IntrinsicData{semantic::Intrinsic::kAtan2, ParamType::kF32, "atan2"}, IntrinsicData{IntrinsicType::kAtan2, ParamType::kF32, "atan2"},
IntrinsicData{semantic::Intrinsic::kCeil, ParamType::kF32, "ceil"}, IntrinsicData{IntrinsicType::kCeil, ParamType::kF32, "ceil"},
IntrinsicData{semantic::Intrinsic::kClamp, ParamType::kF32, "clamp"}, IntrinsicData{IntrinsicType::kClamp, ParamType::kF32, "clamp"},
IntrinsicData{semantic::Intrinsic::kClamp, ParamType::kU32, "clamp"}, IntrinsicData{IntrinsicType::kClamp, ParamType::kU32, "clamp"},
IntrinsicData{semantic::Intrinsic::kCos, ParamType::kF32, "cos"}, IntrinsicData{IntrinsicType::kCos, ParamType::kF32, "cos"},
IntrinsicData{semantic::Intrinsic::kCosh, ParamType::kF32, "cosh"}, IntrinsicData{IntrinsicType::kCosh, ParamType::kF32, "cosh"},
IntrinsicData{semantic::Intrinsic::kCountOneBits, ParamType::kU32, IntrinsicData{IntrinsicType::kCountOneBits, ParamType::kU32,
"countbits"}, "countbits"},
IntrinsicData{semantic::Intrinsic::kCross, ParamType::kF32, "cross"}, IntrinsicData{IntrinsicType::kCross, ParamType::kF32, "cross"},
IntrinsicData{semantic::Intrinsic::kDeterminant, ParamType::kF32, IntrinsicData{IntrinsicType::kDeterminant, ParamType::kF32,
"determinant"}, "determinant"},
IntrinsicData{semantic::Intrinsic::kDistance, ParamType::kF32, IntrinsicData{IntrinsicType::kDistance, ParamType::kF32, "distance"},
"distance"}, IntrinsicData{IntrinsicType::kDot, ParamType::kF32, "dot"},
IntrinsicData{semantic::Intrinsic::kDot, ParamType::kF32, "dot"}, IntrinsicData{IntrinsicType::kDpdx, ParamType::kF32, "ddx"},
IntrinsicData{semantic::Intrinsic::kDpdx, ParamType::kF32, "ddx"}, IntrinsicData{IntrinsicType::kDpdxCoarse, ParamType::kF32,
IntrinsicData{semantic::Intrinsic::kDpdxCoarse, ParamType::kF32,
"ddx_coarse"}, "ddx_coarse"},
IntrinsicData{semantic::Intrinsic::kDpdxFine, ParamType::kF32, IntrinsicData{IntrinsicType::kDpdxFine, ParamType::kF32, "ddx_fine"},
"ddx_fine"}, IntrinsicData{IntrinsicType::kDpdy, ParamType::kF32, "ddy"},
IntrinsicData{semantic::Intrinsic::kDpdy, ParamType::kF32, "ddy"}, IntrinsicData{IntrinsicType::kDpdyCoarse, ParamType::kF32,
IntrinsicData{semantic::Intrinsic::kDpdyCoarse, ParamType::kF32,
"ddy_coarse"}, "ddy_coarse"},
IntrinsicData{semantic::Intrinsic::kDpdyFine, ParamType::kF32, IntrinsicData{IntrinsicType::kDpdyFine, ParamType::kF32, "ddy_fine"},
"ddy_fine"}, IntrinsicData{IntrinsicType::kExp, ParamType::kF32, "exp"},
IntrinsicData{semantic::Intrinsic::kExp, ParamType::kF32, "exp"}, IntrinsicData{IntrinsicType::kExp2, ParamType::kF32, "exp2"},
IntrinsicData{semantic::Intrinsic::kExp2, ParamType::kF32, "exp2"}, IntrinsicData{IntrinsicType::kFaceForward, ParamType::kF32,
IntrinsicData{semantic::Intrinsic::kFaceForward, ParamType::kF32,
"faceforward"}, "faceforward"},
IntrinsicData{semantic::Intrinsic::kFloor, ParamType::kF32, "floor"}, IntrinsicData{IntrinsicType::kFloor, ParamType::kF32, "floor"},
IntrinsicData{semantic::Intrinsic::kFma, ParamType::kF32, "fma"}, IntrinsicData{IntrinsicType::kFma, ParamType::kF32, "fma"},
IntrinsicData{semantic::Intrinsic::kFract, ParamType::kF32, "frac"}, IntrinsicData{IntrinsicType::kFract, ParamType::kF32, "frac"},
IntrinsicData{semantic::Intrinsic::kFwidth, ParamType::kF32, "fwidth"}, IntrinsicData{IntrinsicType::kFwidth, ParamType::kF32, "fwidth"},
IntrinsicData{semantic::Intrinsic::kFwidthCoarse, ParamType::kF32, IntrinsicData{IntrinsicType::kFwidthCoarse, ParamType::kF32, "fwidth"},
"fwidth"}, IntrinsicData{IntrinsicType::kFwidthFine, ParamType::kF32, "fwidth"},
IntrinsicData{semantic::Intrinsic::kFwidthFine, ParamType::kF32, IntrinsicData{IntrinsicType::kInverseSqrt, ParamType::kF32, "rsqrt"},
"fwidth"}, IntrinsicData{IntrinsicType::kIsFinite, ParamType::kF32, "isfinite"},
IntrinsicData{semantic::Intrinsic::kInverseSqrt, ParamType::kF32, IntrinsicData{IntrinsicType::kIsInf, ParamType::kF32, "isinf"},
"rsqrt"}, IntrinsicData{IntrinsicType::kIsNan, ParamType::kF32, "isnan"},
IntrinsicData{semantic::Intrinsic::kIsFinite, ParamType::kF32, IntrinsicData{IntrinsicType::kLdexp, ParamType::kF32, "ldexp"},
"isfinite"}, IntrinsicData{IntrinsicType::kLength, ParamType::kF32, "length"},
IntrinsicData{semantic::Intrinsic::kIsInf, ParamType::kF32, "isinf"}, IntrinsicData{IntrinsicType::kLog, ParamType::kF32, "log"},
IntrinsicData{semantic::Intrinsic::kIsNan, ParamType::kF32, "isnan"}, IntrinsicData{IntrinsicType::kLog2, ParamType::kF32, "log2"},
IntrinsicData{semantic::Intrinsic::kLdexp, ParamType::kF32, "ldexp"}, IntrinsicData{IntrinsicType::kMax, ParamType::kF32, "max"},
IntrinsicData{semantic::Intrinsic::kLength, ParamType::kF32, "length"}, IntrinsicData{IntrinsicType::kMax, ParamType::kU32, "max"},
IntrinsicData{semantic::Intrinsic::kLog, ParamType::kF32, "log"}, IntrinsicData{IntrinsicType::kMin, ParamType::kF32, "min"},
IntrinsicData{semantic::Intrinsic::kLog2, ParamType::kF32, "log2"}, IntrinsicData{IntrinsicType::kMin, ParamType::kU32, "min"},
IntrinsicData{semantic::Intrinsic::kMax, ParamType::kF32, "max"}, IntrinsicData{IntrinsicType::kNormalize, ParamType::kF32, "normalize"},
IntrinsicData{semantic::Intrinsic::kMax, ParamType::kU32, "max"}, IntrinsicData{IntrinsicType::kPow, ParamType::kF32, "pow"},
IntrinsicData{semantic::Intrinsic::kMin, ParamType::kF32, "min"}, IntrinsicData{IntrinsicType::kReflect, ParamType::kF32, "reflect"},
IntrinsicData{semantic::Intrinsic::kMin, ParamType::kU32, "min"}, IntrinsicData{IntrinsicType::kReverseBits, ParamType::kU32,
IntrinsicData{semantic::Intrinsic::kNormalize, ParamType::kF32,
"normalize"},
IntrinsicData{semantic::Intrinsic::kPow, ParamType::kF32, "pow"},
IntrinsicData{semantic::Intrinsic::kReflect, ParamType::kF32,
"reflect"},
IntrinsicData{semantic::Intrinsic::kReverseBits, ParamType::kU32,
"reversebits"}, "reversebits"},
IntrinsicData{semantic::Intrinsic::kRound, ParamType::kU32, "round"}, IntrinsicData{IntrinsicType::kRound, ParamType::kU32, "round"},
IntrinsicData{semantic::Intrinsic::kSign, ParamType::kF32, "sign"}, IntrinsicData{IntrinsicType::kSign, ParamType::kF32, "sign"},
IntrinsicData{semantic::Intrinsic::kSin, ParamType::kF32, "sin"}, IntrinsicData{IntrinsicType::kSin, ParamType::kF32, "sin"},
IntrinsicData{semantic::Intrinsic::kSinh, ParamType::kF32, "sinh"}, IntrinsicData{IntrinsicType::kSinh, ParamType::kF32, "sinh"},
IntrinsicData{semantic::Intrinsic::kSmoothStep, ParamType::kF32, IntrinsicData{IntrinsicType::kSmoothStep, ParamType::kF32,
"smoothstep"}, "smoothstep"},
IntrinsicData{semantic::Intrinsic::kSqrt, ParamType::kF32, "sqrt"}, IntrinsicData{IntrinsicType::kSqrt, ParamType::kF32, "sqrt"},
IntrinsicData{semantic::Intrinsic::kStep, ParamType::kF32, "step"}, IntrinsicData{IntrinsicType::kStep, ParamType::kF32, "step"},
IntrinsicData{semantic::Intrinsic::kTan, ParamType::kF32, "tan"}, IntrinsicData{IntrinsicType::kTan, ParamType::kF32, "tan"},
IntrinsicData{semantic::Intrinsic::kTanh, ParamType::kF32, "tanh"}, IntrinsicData{IntrinsicType::kTanh, ParamType::kF32, "tanh"},
IntrinsicData{semantic::Intrinsic::kTrunc, ParamType::kF32, "trunc"})); IntrinsicData{IntrinsicType::kTrunc, ParamType::kF32, "trunc"}));
TEST_F(HlslGeneratorImplTest_Intrinsic, DISABLED_Intrinsic_IsNormal) { TEST_F(HlslGeneratorImplTest_Intrinsic, DISABLED_Intrinsic_IsNormal) {
FAIL(); FAIL();

148
src/writer/msl/generator_impl.cc
View File

@ -449,7 +449,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) {
return EmitTextureCall(expr, sem); return EmitTextureCall(expr, sem);
} }
if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) { if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) {
if (sem->intrinsic() == semantic::Intrinsic::kPack2x16Float) { if (sem->intrinsic() == semantic::IntrinsicType::kPack2x16Float) {
make_indent(); make_indent();
out_ << "as_type<uint>(half2("; out_ << "as_type<uint>(half2(";
if (!EmitExpression(expr->params()[0])) { if (!EmitExpression(expr->params()[0])) {
@ -580,7 +580,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
TypeOf(params[pidx.texture])->UnwrapAll()->As<type::Texture>(); TypeOf(params[pidx.texture])->UnwrapAll()->As<type::Texture>();
switch (sem->intrinsic()) { switch (sem->intrinsic()) {
case semantic::Intrinsic::kTextureDimensions: { case semantic::IntrinsicType::kTextureDimensions: {
std::vector<const char*> dims; std::vector<const char*> dims;
switch (texture_type->dim()) { switch (texture_type->dim()) {
case type::TextureDimension::kNone: case type::TextureDimension::kNone:
@ -634,7 +634,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
} }
return true; return true;
} }
case semantic::Intrinsic::kTextureNumLayers: { case semantic::IntrinsicType::kTextureNumLayers: {
out_ << "int("; out_ << "int(";
if (!EmitExpression(params[pidx.texture])) { if (!EmitExpression(params[pidx.texture])) {
return false; return false;
@ -642,7 +642,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
out_ << ".get_array_size())"; out_ << ".get_array_size())";
return true; return true;
} }
case semantic::Intrinsic::kTextureNumLevels: { case semantic::IntrinsicType::kTextureNumLevels: {
out_ << "int("; out_ << "int(";
if (!EmitExpression(params[pidx.texture])) { if (!EmitExpression(params[pidx.texture])) {
return false; return false;
@ -650,7 +650,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
out_ << ".get_num_mip_levels())"; out_ << ".get_num_mip_levels())";
return true; return true;
} }
case semantic::Intrinsic::kTextureNumSamples: { case semantic::IntrinsicType::kTextureNumSamples: {
out_ << "int("; out_ << "int(";
if (!EmitExpression(params[pidx.texture])) { if (!EmitExpression(params[pidx.texture])) {
return false; return false;
@ -668,20 +668,20 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr,
bool lod_param_is_named = true; bool lod_param_is_named = true;
switch (sem->intrinsic()) { switch (sem->intrinsic()) {
case semantic::Intrinsic::kTextureSample: case semantic::IntrinsicType::kTextureSample:
case semantic::Intrinsic::kTextureSampleBias: case semantic::IntrinsicType::kTextureSampleBias:
case semantic::Intrinsic::kTextureSampleLevel: case semantic::IntrinsicType::kTextureSampleLevel:
case semantic::Intrinsic::kTextureSampleGrad: case semantic::IntrinsicType::kTextureSampleGrad:
out_ << ".sample("; out_ << ".sample(";
break; break;
case semantic::Intrinsic::kTextureSampleCompare: case semantic::IntrinsicType::kTextureSampleCompare:
out_ << ".sample_compare("; out_ << ".sample_compare(";
break; break;
case semantic::Intrinsic::kTextureLoad: case semantic::IntrinsicType::kTextureLoad:
out_ << ".read("; out_ << ".read(";
lod_param_is_named = false; lod_param_is_named = false;
break; break;
case semantic::Intrinsic::kTextureStore: case semantic::IntrinsicType::kTextureStore:
out_ << ".write("; out_ << ".write(";
break; break;
default: default:
@ -780,118 +780,118 @@ std::string GeneratorImpl::generate_builtin_name(
const semantic::IntrinsicCall* call) { const semantic::IntrinsicCall* call) {
std::string out = "metal::"; std::string out = "metal::";
switch (call->intrinsic()) { switch (call->intrinsic()) {
case semantic::Intrinsic::kAcos: case semantic::IntrinsicType::kAcos:
case semantic::Intrinsic::kAll: case semantic::IntrinsicType::kAll:
case semantic::Intrinsic::kAny: case semantic::IntrinsicType::kAny:
case semantic::Intrinsic::kAsin: case semantic::IntrinsicType::kAsin:
case semantic::Intrinsic::kAtan: case semantic::IntrinsicType::kAtan:
case semantic::Intrinsic::kAtan2: case semantic::IntrinsicType::kAtan2:
case semantic::Intrinsic::kCeil: case semantic::IntrinsicType::kCeil:
case semantic::Intrinsic::kCos: case semantic::IntrinsicType::kCos:
case semantic::Intrinsic::kCosh: case semantic::IntrinsicType::kCosh:
case semantic::Intrinsic::kCross: case semantic::IntrinsicType::kCross:
case semantic::Intrinsic::kDeterminant: case semantic::IntrinsicType::kDeterminant:
case semantic::Intrinsic::kDistance: case semantic::IntrinsicType::kDistance:
case semantic::Intrinsic::kDot: case semantic::IntrinsicType::kDot:
case semantic::Intrinsic::kExp: case semantic::IntrinsicType::kExp:
case semantic::Intrinsic::kExp2: case semantic::IntrinsicType::kExp2:
case semantic::Intrinsic::kFloor: case semantic::IntrinsicType::kFloor:
case semantic::Intrinsic::kFma: case semantic::IntrinsicType::kFma:
case semantic::Intrinsic::kFract: case semantic::IntrinsicType::kFract:
case semantic::Intrinsic::kLength: case semantic::IntrinsicType::kLength:
case semantic::Intrinsic::kLdexp: case semantic::IntrinsicType::kLdexp:
case semantic::Intrinsic::kLog: case semantic::IntrinsicType::kLog:
case semantic::Intrinsic::kLog2: case semantic::IntrinsicType::kLog2:
case semantic::Intrinsic::kMix: case semantic::IntrinsicType::kMix:
case semantic::Intrinsic::kNormalize: case semantic::IntrinsicType::kNormalize:
case semantic::Intrinsic::kPow: case semantic::IntrinsicType::kPow:
case semantic::Intrinsic::kReflect: case semantic::IntrinsicType::kReflect:
case semantic::Intrinsic::kRound: case semantic::IntrinsicType::kRound:
case semantic::Intrinsic::kSelect: case semantic::IntrinsicType::kSelect:
case semantic::Intrinsic::kSin: case semantic::IntrinsicType::kSin:
case semantic::Intrinsic::kSinh: case semantic::IntrinsicType::kSinh:
case semantic::Intrinsic::kSqrt: case semantic::IntrinsicType::kSqrt:
case semantic::Intrinsic::kStep: case semantic::IntrinsicType::kStep:
case semantic::Intrinsic::kTan: case semantic::IntrinsicType::kTan:
case semantic::Intrinsic::kTanh: case semantic::IntrinsicType::kTanh:
case semantic::Intrinsic::kTrunc: case semantic::IntrinsicType::kTrunc:
case semantic::Intrinsic::kSign: case semantic::IntrinsicType::kSign:
case semantic::Intrinsic::kClamp: case semantic::IntrinsicType::kClamp:
out += semantic::intrinsic::str(call->intrinsic()); out += semantic::intrinsic::str(call->intrinsic());
break; break;
case semantic::Intrinsic::kAbs: case semantic::IntrinsicType::kAbs:
if (call->Type()->is_float_scalar_or_vector()) { if (call->Type()->is_float_scalar_or_vector()) {
out += "fabs"; out += "fabs";
} else { } else {
out += "abs"; out += "abs";
} }
break; break;
case semantic::Intrinsic::kCountOneBits: case semantic::IntrinsicType::kCountOneBits:
out += "popcount"; out += "popcount";
break; break;
case semantic::Intrinsic::kDpdx: case semantic::IntrinsicType::kDpdx:
case semantic::Intrinsic::kDpdxCoarse: case semantic::IntrinsicType::kDpdxCoarse:
case semantic::Intrinsic::kDpdxFine: case semantic::IntrinsicType::kDpdxFine:
out += "dfdx"; out += "dfdx";
break; break;
case semantic::Intrinsic::kDpdy: case semantic::IntrinsicType::kDpdy:
case semantic::Intrinsic::kDpdyCoarse: case semantic::IntrinsicType::kDpdyCoarse:
case semantic::Intrinsic::kDpdyFine: case semantic::IntrinsicType::kDpdyFine:
out += "dfdy"; out += "dfdy";
break; break;
case semantic::Intrinsic::kFwidth: case semantic::IntrinsicType::kFwidth:
case semantic::Intrinsic::kFwidthCoarse: case semantic::IntrinsicType::kFwidthCoarse:
case semantic::Intrinsic::kFwidthFine: case semantic::IntrinsicType::kFwidthFine:
out += "fwidth"; out += "fwidth";
break; break;
case semantic::Intrinsic::kIsFinite: case semantic::IntrinsicType::kIsFinite:
out += "isfinite"; out += "isfinite";
break; break;
case semantic::Intrinsic::kIsInf: case semantic::IntrinsicType::kIsInf:
out += "isinf"; out += "isinf";
break; break;
case semantic::Intrinsic::kIsNan: case semantic::IntrinsicType::kIsNan:
out += "isnan"; out += "isnan";
break; break;
case semantic::Intrinsic::kIsNormal: case semantic::IntrinsicType::kIsNormal:
out += "isnormal"; out += "isnormal";
break; break;
case semantic::Intrinsic::kMax: case semantic::IntrinsicType::kMax:
if (call->Type()->is_float_scalar_or_vector()) { if (call->Type()->is_float_scalar_or_vector()) {
out += "fmax"; out += "fmax";
} else { } else {
out += "max"; out += "max";
} }
break; break;
case semantic::Intrinsic::kMin: case semantic::IntrinsicType::kMin:
if (call->Type()->is_float_scalar_or_vector()) { if (call->Type()->is_float_scalar_or_vector()) {
out += "fmin"; out += "fmin";
} else { } else {
out += "min"; out += "min";
} }
break; break;
case semantic::Intrinsic::kFaceForward: case semantic::IntrinsicType::kFaceForward:
out += "faceforward"; out += "faceforward";
break; break;
case semantic::Intrinsic::kPack4x8Snorm: case semantic::IntrinsicType::kPack4x8Snorm:
out += "pack_float_to_snorm4x8"; out += "pack_float_to_snorm4x8";
break; break;
case semantic::Intrinsic::kPack4x8Unorm: case semantic::IntrinsicType::kPack4x8Unorm:
out += "pack_float_to_unorm4x8"; out += "pack_float_to_unorm4x8";
break; break;
case semantic::Intrinsic::kPack2x16Snorm: case semantic::IntrinsicType::kPack2x16Snorm:
out += "pack_float_to_snorm2x16"; out += "pack_float_to_snorm2x16";
break; break;
case semantic::Intrinsic::kPack2x16Unorm: case semantic::IntrinsicType::kPack2x16Unorm:
out += "pack_float_to_unorm2x16"; out += "pack_float_to_unorm2x16";
break; break;
case semantic::Intrinsic::kReverseBits: case semantic::IntrinsicType::kReverseBits:
out += "reverse_bits"; out += "reverse_bits";
break; break;
case semantic::Intrinsic::kSmoothStep: case semantic::IntrinsicType::kSmoothStep:
out += "smoothstep"; out += "smoothstep";
break; break;
case semantic::Intrinsic::kInverseSqrt: case semantic::IntrinsicType::kInverseSqrt:
out += "rsqrt"; out += "rsqrt";
break; break;
default: default:

295
src/writer/msl/generator_impl_intrinsic_test.cc
View File

@ -30,6 +30,8 @@ namespace writer {
namespace msl { namespace msl {
namespace { namespace {
using IntrinsicType = semantic::IntrinsicType;
using MslGeneratorImplTest = TestHelper; using MslGeneratorImplTest = TestHelper;
enum class ParamType { enum class ParamType {
@ -39,7 +41,7 @@ enum class ParamType {
}; };
struct IntrinsicData { struct IntrinsicData {
semantic::Intrinsic intrinsic; IntrinsicType intrinsic;
ParamType type; ParamType type;
const char* msl_name; const char* msl_name;
}; };
@ -60,97 +62,97 @@ inline std::ostream& operator<<(std::ostream& out, IntrinsicData data) {
return out; return out;
} }
ast::CallExpression* GenerateCall(semantic::Intrinsic intrinsic, ast::CallExpression* GenerateCall(IntrinsicType intrinsic,
ParamType type, ParamType type,
ProgramBuilder* builder) { ProgramBuilder* builder) {
std::string name; std::string name;
std::ostringstream str(name); std::ostringstream str(name);
str << intrinsic; str << intrinsic;
switch (intrinsic) { switch (intrinsic) {
case semantic::Intrinsic::kAcos: case IntrinsicType::kAcos:
case semantic::Intrinsic::kAsin: case IntrinsicType::kAsin:
case semantic::Intrinsic::kAtan: case IntrinsicType::kAtan:
case semantic::Intrinsic::kCeil: case IntrinsicType::kCeil:
case semantic::Intrinsic::kCos: case IntrinsicType::kCos:
case semantic::Intrinsic::kCosh: case IntrinsicType::kCosh:
case semantic::Intrinsic::kDpdx: case IntrinsicType::kDpdx:
case semantic::Intrinsic::kDpdxCoarse: case IntrinsicType::kDpdxCoarse:
case semantic::Intrinsic::kDpdxFine: case IntrinsicType::kDpdxFine:
case semantic::Intrinsic::kDpdy: case IntrinsicType::kDpdy:
case semantic::Intrinsic::kDpdyCoarse: case IntrinsicType::kDpdyCoarse:
case semantic::Intrinsic::kDpdyFine: case IntrinsicType::kDpdyFine:
case semantic::Intrinsic::kExp: case IntrinsicType::kExp:
case semantic::Intrinsic::kExp2: case IntrinsicType::kExp2:
case semantic::Intrinsic::kFloor: case IntrinsicType::kFloor:
case semantic::Intrinsic::kFract: case IntrinsicType::kFract:
case semantic::Intrinsic::kFwidth: case IntrinsicType::kFwidth:
case semantic::Intrinsic::kFwidthCoarse: case IntrinsicType::kFwidthCoarse:
case semantic::Intrinsic::kFwidthFine: case IntrinsicType::kFwidthFine:
case semantic::Intrinsic::kInverseSqrt: case IntrinsicType::kInverseSqrt:
case semantic::Intrinsic::kIsFinite: case IntrinsicType::kIsFinite:
case semantic::Intrinsic::kIsInf: case IntrinsicType::kIsInf:
case semantic::Intrinsic::kIsNan: case IntrinsicType::kIsNan:
case semantic::Intrinsic::kIsNormal: case IntrinsicType::kIsNormal:
case semantic::Intrinsic::kLdexp: case IntrinsicType::kLdexp:
case semantic::Intrinsic::kLength: case IntrinsicType::kLength:
case semantic::Intrinsic::kLog: case IntrinsicType::kLog:
case semantic::Intrinsic::kLog2: case IntrinsicType::kLog2:
case semantic::Intrinsic::kNormalize: case IntrinsicType::kNormalize:
case semantic::Intrinsic::kPack4x8Snorm: case IntrinsicType::kPack4x8Snorm:
case semantic::Intrinsic::kPack4x8Unorm: case IntrinsicType::kPack4x8Unorm:
case semantic::Intrinsic::kReflect: case IntrinsicType::kReflect:
case semantic::Intrinsic::kRound: case IntrinsicType::kRound:
case semantic::Intrinsic::kSin: case IntrinsicType::kSin:
case semantic::Intrinsic::kSinh: case IntrinsicType::kSinh:
case semantic::Intrinsic::kSqrt: case IntrinsicType::kSqrt:
case semantic::Intrinsic::kTan: case IntrinsicType::kTan:
case semantic::Intrinsic::kTanh: case IntrinsicType::kTanh:
case semantic::Intrinsic::kTrunc: case IntrinsicType::kTrunc:
case semantic::Intrinsic::kSign: case IntrinsicType::kSign:
return builder->Call(str.str(), "f1"); return builder->Call(str.str(), "f1");
case semantic::Intrinsic::kAtan2: case IntrinsicType::kAtan2:
case semantic::Intrinsic::kCross: case IntrinsicType::kCross:
case semantic::Intrinsic::kDot: case IntrinsicType::kDot:
case semantic::Intrinsic::kDistance: case IntrinsicType::kDistance:
case semantic::Intrinsic::kPow: case IntrinsicType::kPow:
case semantic::Intrinsic::kStep: case IntrinsicType::kStep:
return builder->Call(str.str(), "f1", "f2"); return builder->Call(str.str(), "f1", "f2");
case semantic::Intrinsic::kFma: case IntrinsicType::kFma:
case semantic::Intrinsic::kMix: case IntrinsicType::kMix:
case semantic::Intrinsic::kFaceForward: case IntrinsicType::kFaceForward:
case semantic::Intrinsic::kSmoothStep: case IntrinsicType::kSmoothStep:
return builder->Call(str.str(), "f1", "f2", "f3"); return builder->Call(str.str(), "f1", "f2", "f3");
case semantic::Intrinsic::kAll: case IntrinsicType::kAll:
case semantic::Intrinsic::kAny: case IntrinsicType::kAny:
return builder->Call(str.str(), "b1"); return builder->Call(str.str(), "b1");
case semantic::Intrinsic::kAbs: case IntrinsicType::kAbs:
if (type == ParamType::kF32) { if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1"); return builder->Call(str.str(), "f1");
} else { } else {
return builder->Call(str.str(), "u1"); return builder->Call(str.str(), "u1");
} }
case semantic::Intrinsic::kCountOneBits: case IntrinsicType::kCountOneBits:
case semantic::Intrinsic::kReverseBits: case IntrinsicType::kReverseBits:
return builder->Call(str.str(), "u1"); return builder->Call(str.str(), "u1");
case semantic::Intrinsic::kMax: case IntrinsicType::kMax:
case semantic::Intrinsic::kMin: case IntrinsicType::kMin:
if (type == ParamType::kF32) { if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1", "f2"); return builder->Call(str.str(), "f1", "f2");
} else { } else {
return builder->Call(str.str(), "u1", "u2"); return builder->Call(str.str(), "u1", "u2");
} }
case semantic::Intrinsic::kClamp: case IntrinsicType::kClamp:
if (type == ParamType::kF32) { if (type == ParamType::kF32) {
return builder->Call(str.str(), "f1", "f2", "f3"); return builder->Call(str.str(), "f1", "f2", "f3");
} else { } else {
return builder->Call(str.str(), "u1", "u2", "u3"); return builder->Call(str.str(), "u1", "u2", "u3");
} }
case semantic::Intrinsic::kSelect: case IntrinsicType::kSelect:
return builder->Call(str.str(), "f1", "f2", "b1"); return builder->Call(str.str(), "f1", "f2", "b1");
case semantic::Intrinsic::kDeterminant: case IntrinsicType::kDeterminant:
return builder->Call(str.str(), "m1"); return builder->Call(str.str(), "m1");
case semantic::Intrinsic::kPack2x16Snorm: case IntrinsicType::kPack2x16Snorm:
case semantic::Intrinsic::kPack2x16Unorm: case IntrinsicType::kPack2x16Unorm:
return builder->Call(str.str(), "f4"); return builder->Call(str.str(), "f4");
default: default:
break; break;
@ -189,125 +191,90 @@ INSTANTIATE_TEST_SUITE_P(
MslGeneratorImplTest, MslGeneratorImplTest,
MslIntrinsicTest, MslIntrinsicTest,
testing::Values( testing::Values(
IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kF32, IntrinsicData{IntrinsicType::kAbs, ParamType::kF32, "metal::fabs"},
"metal::fabs"}, IntrinsicData{IntrinsicType::kAbs, ParamType::kU32, "metal::abs"},
IntrinsicData{semantic::Intrinsic::kAbs, ParamType::kU32, "metal::abs"}, IntrinsicData{IntrinsicType::kAcos, ParamType::kF32, "metal::acos"},
IntrinsicData{semantic::Intrinsic::kAcos, ParamType::kF32, IntrinsicData{IntrinsicType::kAll, ParamType::kBool, "metal::all"},
"metal::acos"}, IntrinsicData{IntrinsicType::kAny, ParamType::kBool, "metal::any"},
IntrinsicData{semantic::Intrinsic::kAll, ParamType::kBool, IntrinsicData{IntrinsicType::kAsin, ParamType::kF32, "metal::asin"},
"metal::all"}, IntrinsicData{IntrinsicType::kAtan, ParamType::kF32, "metal::atan"},
IntrinsicData{semantic::Intrinsic::kAny, ParamType::kBool, IntrinsicData{IntrinsicType::kAtan2, ParamType::kF32, "metal::atan2"},
"metal::any"}, IntrinsicData{IntrinsicType::kCeil, ParamType::kF32, "metal::ceil"},
IntrinsicData{semantic::Intrinsic::kAsin, ParamType::kF32, IntrinsicData{IntrinsicType::kClamp, ParamType::kF32, "metal::clamp"},
"metal::asin"}, IntrinsicData{IntrinsicType::kClamp, ParamType::kU32, "metal::clamp"},
IntrinsicData{semantic::Intrinsic::kAtan, ParamType::kF32, IntrinsicData{IntrinsicType::kCos, ParamType::kF32, "metal::cos"},
"metal::atan"}, IntrinsicData{IntrinsicType::kCosh, ParamType::kF32, "metal::cosh"},
IntrinsicData{semantic::Intrinsic::kAtan2, ParamType::kF32, IntrinsicData{IntrinsicType::kCountOneBits, ParamType::kU32,
"metal::atan2"},
IntrinsicData{semantic::Intrinsic::kCeil, ParamType::kF32,
"metal::ceil"},
IntrinsicData{semantic::Intrinsic::kClamp, ParamType::kF32,
"metal::clamp"},
IntrinsicData{semantic::Intrinsic::kClamp, ParamType::kU32,
"metal::clamp"},
IntrinsicData{semantic::Intrinsic::kCos, ParamType::kF32, "metal::cos"},
IntrinsicData{semantic::Intrinsic::kCosh, ParamType::kF32,
"metal::cosh"},
IntrinsicData{semantic::Intrinsic::kCountOneBits, ParamType::kU32,
"metal::popcount"}, "metal::popcount"},
IntrinsicData{semantic::Intrinsic::kCross, ParamType::kF32, IntrinsicData{IntrinsicType::kCross, ParamType::kF32, "metal::cross"},
"metal::cross"}, IntrinsicData{IntrinsicType::kDeterminant, ParamType::kF32,
IntrinsicData{semantic::Intrinsic::kDeterminant, ParamType::kF32,
"metal::determinant"}, "metal::determinant"},
IntrinsicData{semantic::Intrinsic::kDistance, ParamType::kF32, IntrinsicData{IntrinsicType::kDistance, ParamType::kF32,
"metal::distance"}, "metal::distance"},
IntrinsicData{semantic::Intrinsic::kDot, ParamType::kF32, "metal::dot"}, IntrinsicData{IntrinsicType::kDot, ParamType::kF32, "metal::dot"},
IntrinsicData{semantic::Intrinsic::kDpdx, ParamType::kF32, IntrinsicData{IntrinsicType::kDpdx, ParamType::kF32, "metal::dfdx"},
IntrinsicData{IntrinsicType::kDpdxCoarse, ParamType::kF32,
"metal::dfdx"}, "metal::dfdx"},
IntrinsicData{semantic::Intrinsic::kDpdxCoarse, ParamType::kF32, IntrinsicData{IntrinsicType::kDpdxFine, ParamType::kF32, "metal::dfdx"},
"metal::dfdx"}, IntrinsicData{IntrinsicType::kDpdy, ParamType::kF32, "metal::dfdy"},
IntrinsicData{semantic::Intrinsic::kDpdxFine, ParamType::kF32, IntrinsicData{IntrinsicType::kDpdyCoarse, ParamType::kF32,
"metal::dfdx"},
IntrinsicData{semantic::Intrinsic::kDpdy, ParamType::kF32,
"metal::dfdy"}, "metal::dfdy"},
IntrinsicData{semantic::Intrinsic::kDpdyCoarse, ParamType::kF32, IntrinsicData{IntrinsicType::kDpdyFine, ParamType::kF32, "metal::dfdy"},
"metal::dfdy"}, IntrinsicData{IntrinsicType::kExp, ParamType::kF32, "metal::exp"},
IntrinsicData{semantic::Intrinsic::kDpdyFine, ParamType::kF32, IntrinsicData{IntrinsicType::kExp2, ParamType::kF32, "metal::exp2"},
"metal::dfdy"}, IntrinsicData{IntrinsicType::kFaceForward, ParamType::kF32,
IntrinsicData{semantic::Intrinsic::kExp, ParamType::kF32, "metal::exp"},
IntrinsicData{semantic::Intrinsic::kExp2, ParamType::kF32,
"metal::exp2"},
IntrinsicData{semantic::Intrinsic::kFaceForward, ParamType::kF32,
"metal::faceforward"}, "metal::faceforward"},
IntrinsicData{semantic::Intrinsic::kFloor, ParamType::kF32, IntrinsicData{IntrinsicType::kFloor, ParamType::kF32, "metal::floor"},
"metal::floor"}, IntrinsicData{IntrinsicType::kFma, ParamType::kF32, "metal::fma"},
IntrinsicData{semantic::Intrinsic::kFma, ParamType::kF32, "metal::fma"}, IntrinsicData{IntrinsicType::kFract, ParamType::kF32, "metal::fract"},
IntrinsicData{semantic::Intrinsic::kFract, ParamType::kF32, IntrinsicData{IntrinsicType::kFwidth, ParamType::kF32, "metal::fwidth"},
"metal::fract"}, IntrinsicData{IntrinsicType::kFwidthCoarse, ParamType::kF32,
IntrinsicData{semantic::Intrinsic::kFwidth, ParamType::kF32,
"metal::fwidth"}, "metal::fwidth"},
IntrinsicData{semantic::Intrinsic::kFwidthCoarse, ParamType::kF32, IntrinsicData{IntrinsicType::kFwidthFine, ParamType::kF32,
"metal::fwidth"}, "metal::fwidth"},
IntrinsicData{semantic::Intrinsic::kFwidthFine, ParamType::kF32, IntrinsicData{IntrinsicType::kInverseSqrt, ParamType::kF32,
"metal::fwidth"},
IntrinsicData{semantic::Intrinsic::kInverseSqrt, ParamType::kF32,
"metal::rsqrt"}, "metal::rsqrt"},
IntrinsicData{semantic::Intrinsic::kIsFinite, ParamType::kF32, IntrinsicData{IntrinsicType::kIsFinite, ParamType::kF32,
"metal::isfinite"}, "metal::isfinite"},
IntrinsicData{semantic::Intrinsic::kIsInf, ParamType::kF32, IntrinsicData{IntrinsicType::kIsInf, ParamType::kF32, "metal::isinf"},
"metal::isinf"}, IntrinsicData{IntrinsicType::kIsNan, ParamType::kF32, "metal::isnan"},
IntrinsicData{semantic::Intrinsic::kIsNan, ParamType::kF32, IntrinsicData{IntrinsicType::kIsNormal, ParamType::kF32,
"metal::isnan"},
IntrinsicData{semantic::Intrinsic::kIsNormal, ParamType::kF32,
"metal::isnormal"}, "metal::isnormal"},
IntrinsicData{semantic::Intrinsic::kLdexp, ParamType::kF32, IntrinsicData{IntrinsicType::kLdexp, ParamType::kF32, "metal::ldexp"},
"metal::ldexp"}, IntrinsicData{IntrinsicType::kLength, ParamType::kF32, "metal::length"},
IntrinsicData{semantic::Intrinsic::kLength, ParamType::kF32, IntrinsicData{IntrinsicType::kLog, ParamType::kF32, "metal::log"},
"metal::length"}, IntrinsicData{IntrinsicType::kLog2, ParamType::kF32, "metal::log2"},
IntrinsicData{semantic::Intrinsic::kLog, ParamType::kF32, "metal::log"}, IntrinsicData{IntrinsicType::kMax, ParamType::kF32, "metal::fmax"},
IntrinsicData{semantic::Intrinsic::kLog2, ParamType::kF32, IntrinsicData{IntrinsicType::kMax, ParamType::kU32, "metal::max"},
"metal::log2"}, IntrinsicData{IntrinsicType::kMin, ParamType::kF32, "metal::fmin"},
IntrinsicData{semantic::Intrinsic::kMax, ParamType::kF32, IntrinsicData{IntrinsicType::kMin, ParamType::kU32, "metal::min"},
"metal::fmax"}, IntrinsicData{IntrinsicType::kNormalize, ParamType::kF32,
IntrinsicData{semantic::Intrinsic::kMax, ParamType::kU32, "metal::max"},
IntrinsicData{semantic::Intrinsic::kMin, ParamType::kF32,
"metal::fmin"},
IntrinsicData{semantic::Intrinsic::kMin, ParamType::kU32, "metal::min"},
IntrinsicData{semantic::Intrinsic::kNormalize, ParamType::kF32,
"metal::normalize"}, "metal::normalize"},
IntrinsicData{semantic::Intrinsic::kPack4x8Snorm, ParamType::kF32, IntrinsicData{IntrinsicType::kPack4x8Snorm, ParamType::kF32,
"metal::pack_float_to_snorm4x8"}, "metal::pack_float_to_snorm4x8"},
IntrinsicData{semantic::Intrinsic::kPack4x8Unorm, ParamType::kF32, IntrinsicData{IntrinsicType::kPack4x8Unorm, ParamType::kF32,
"metal::pack_float_to_unorm4x8"}, "metal::pack_float_to_unorm4x8"},
IntrinsicData{semantic::Intrinsic::kPack2x16Snorm, ParamType::kF32, IntrinsicData{IntrinsicType::kPack2x16Snorm, ParamType::kF32,
"metal::pack_float_to_snorm2x16"}, "metal::pack_float_to_snorm2x16"},
IntrinsicData{semantic::Intrinsic::kPack2x16Unorm, ParamType::kF32, IntrinsicData{IntrinsicType::kPack2x16Unorm, ParamType::kF32,
"metal::pack_float_to_unorm2x16"}, "metal::pack_float_to_unorm2x16"},
IntrinsicData{semantic::Intrinsic::kPow, ParamType::kF32, "metal::pow"}, IntrinsicData{IntrinsicType::kPow, ParamType::kF32, "metal::pow"},
IntrinsicData{semantic::Intrinsic::kReflect, ParamType::kF32, IntrinsicData{IntrinsicType::kReflect, ParamType::kF32,
"metal::reflect"}, "metal::reflect"},
IntrinsicData{semantic::Intrinsic::kReverseBits, ParamType::kU32, IntrinsicData{IntrinsicType::kReverseBits, ParamType::kU32,
"metal::reverse_bits"}, "metal::reverse_bits"},
IntrinsicData{semantic::Intrinsic::kRound, ParamType::kU32, IntrinsicData{IntrinsicType::kRound, ParamType::kU32, "metal::round"},
"metal::round"}, IntrinsicData{IntrinsicType::kSelect, ParamType::kF32, "metal::select"},
IntrinsicData{semantic::Intrinsic::kSelect, ParamType::kF32, IntrinsicData{IntrinsicType::kSign, ParamType::kF32, "metal::sign"},
"metal::select"}, IntrinsicData{IntrinsicType::kSin, ParamType::kF32, "metal::sin"},
IntrinsicData{semantic::Intrinsic::kSign, ParamType::kF32, IntrinsicData{IntrinsicType::kSinh, ParamType::kF32, "metal::sinh"},
"metal::sign"}, IntrinsicData{IntrinsicType::kSmoothStep, ParamType::kF32,
IntrinsicData{semantic::Intrinsic::kSin, ParamType::kF32, "metal::sin"},
IntrinsicData{semantic::Intrinsic::kSinh, ParamType::kF32,
"metal::sinh"},
IntrinsicData{semantic::Intrinsic::kSmoothStep, ParamType::kF32,
"metal::smoothstep"}, "metal::smoothstep"},
IntrinsicData{semantic::Intrinsic::kSqrt, ParamType::kF32, IntrinsicData{IntrinsicType::kSqrt, ParamType::kF32, "metal::sqrt"},
"metal::sqrt"}, IntrinsicData{IntrinsicType::kStep, ParamType::kF32, "metal::step"},
IntrinsicData{semantic::Intrinsic::kStep, ParamType::kF32, IntrinsicData{IntrinsicType::kTan, ParamType::kF32, "metal::tan"},
"metal::step"}, IntrinsicData{IntrinsicType::kTanh, ParamType::kF32, "metal::tanh"},
IntrinsicData{semantic::Intrinsic::kTan, ParamType::kF32, "metal::tan"}, IntrinsicData{IntrinsicType::kTrunc, ParamType::kF32, "metal::trunc"}));
IntrinsicData{semantic::Intrinsic::kTanh, ParamType::kF32,
"metal::tanh"},
IntrinsicData{semantic::Intrinsic::kTrunc, ParamType::kF32,
"metal::trunc"}));
TEST_F(MslGeneratorImplTest, Intrinsic_Call) { TEST_F(MslGeneratorImplTest, Intrinsic_Call) {
Global("param1", ast::StorageClass::kFunction, ty.vec2<f32>()); Global("param1", ast::StorageClass::kFunction, ty.vec2<f32>());

155
src/writer/spirv/builder.cc
View File

@ -87,6 +87,8 @@ namespace writer {
namespace spirv { namespace spirv {
namespace { namespace {
using IntrinsicType = semantic::IntrinsicType;
const char kGLSLstd450[] = "GLSL.std.450"; const char kGLSLstd450[] = "GLSL.std.450";
uint32_t size_of(const InstructionList& instructions) { uint32_t size_of(const InstructionList& instructions) {
@ -166,26 +168,25 @@ type::Matrix* GetNestedMatrixType(type::Type* type) {
return type->As<type::Matrix>(); return type->As<type::Matrix>();
} }
uint32_t intrinsic_to_glsl_method(type::Type* type, uint32_t intrinsic_to_glsl_method(type::Type* type, IntrinsicType intrinsic) {
semantic::Intrinsic intrinsic) {
switch (intrinsic) { switch (intrinsic) {
case semantic::Intrinsic::kAbs: case IntrinsicType::kAbs:
if (type->is_float_scalar_or_vector()) { if (type->is_float_scalar_or_vector()) {
return GLSLstd450FAbs; return GLSLstd450FAbs;
} else { } else {
return GLSLstd450SAbs; return GLSLstd450SAbs;
} }
case semantic::Intrinsic::kAcos: case IntrinsicType::kAcos:
return GLSLstd450Acos; return GLSLstd450Acos;
case semantic::Intrinsic::kAsin: case IntrinsicType::kAsin:
return GLSLstd450Asin; return GLSLstd450Asin;
case semantic::Intrinsic::kAtan: case IntrinsicType::kAtan:
return GLSLstd450Atan; return GLSLstd450Atan;
case semantic::Intrinsic::kAtan2: case IntrinsicType::kAtan2:
return GLSLstd450Atan2; return GLSLstd450Atan2;
case semantic::Intrinsic::kCeil: case IntrinsicType::kCeil:
return GLSLstd450Ceil; return GLSLstd450Ceil;
case semantic::Intrinsic::kClamp: case IntrinsicType::kClamp:
if (type->is_float_scalar_or_vector()) { if (type->is_float_scalar_or_vector()) {
return GLSLstd450NClamp; return GLSLstd450NClamp;
} else if (type->is_unsigned_scalar_or_vector()) { } else if (type->is_unsigned_scalar_or_vector()) {
@ -193,41 +194,41 @@ uint32_t intrinsic_to_glsl_method(type::Type* type,
} else { } else {
return GLSLstd450SClamp; return GLSLstd450SClamp;
} }
case semantic::Intrinsic::kCos: case IntrinsicType::kCos:
return GLSLstd450Cos; return GLSLstd450Cos;
case semantic::Intrinsic::kCosh: case IntrinsicType::kCosh:
return GLSLstd450Cosh; return GLSLstd450Cosh;
case semantic::Intrinsic::kCross: case IntrinsicType::kCross:
return GLSLstd450Cross; return GLSLstd450Cross;
case semantic::Intrinsic::kDeterminant: case IntrinsicType::kDeterminant:
return GLSLstd450Determinant; return GLSLstd450Determinant;
case semantic::Intrinsic::kDistance: case IntrinsicType::kDistance:
return GLSLstd450Distance; return GLSLstd450Distance;
case semantic::Intrinsic::kExp: case IntrinsicType::kExp:
return GLSLstd450Exp; return GLSLstd450Exp;
case semantic::Intrinsic::kExp2: case IntrinsicType::kExp2:
return GLSLstd450Exp2; return GLSLstd450Exp2;
case semantic::Intrinsic::kFaceForward: case IntrinsicType::kFaceForward:
return GLSLstd450FaceForward; return GLSLstd450FaceForward;
case semantic::Intrinsic::kFloor: case IntrinsicType::kFloor:
return GLSLstd450Floor; return GLSLstd450Floor;
case semantic::Intrinsic::kFma: case IntrinsicType::kFma:
return GLSLstd450Fma; return GLSLstd450Fma;
case semantic::Intrinsic::kFract: case IntrinsicType::kFract:
return GLSLstd450Fract; return GLSLstd450Fract;
case semantic::Intrinsic::kFrexp: case IntrinsicType::kFrexp:
return GLSLstd450Frexp; return GLSLstd450Frexp;
case semantic::Intrinsic::kInverseSqrt: case IntrinsicType::kInverseSqrt:
return GLSLstd450InverseSqrt; return GLSLstd450InverseSqrt;
case semantic::Intrinsic::kLdexp: case IntrinsicType::kLdexp:
return GLSLstd450Ldexp; return GLSLstd450Ldexp;
case semantic::Intrinsic::kLength: case IntrinsicType::kLength:
return GLSLstd450Length; return GLSLstd450Length;
case semantic::Intrinsic::kLog: case IntrinsicType::kLog:
return GLSLstd450Log; return GLSLstd450Log;
case semantic::Intrinsic::kLog2: case IntrinsicType::kLog2:
return GLSLstd450Log2; return GLSLstd450Log2;
case semantic::Intrinsic::kMax: case IntrinsicType::kMax:
if (type->is_float_scalar_or_vector()) { if (type->is_float_scalar_or_vector()) {
return GLSLstd450NMax; return GLSLstd450NMax;
} else if (type->is_unsigned_scalar_or_vector()) { } else if (type->is_unsigned_scalar_or_vector()) {
@ -235,7 +236,7 @@ uint32_t intrinsic_to_glsl_method(type::Type* type,
} else { } else {
return GLSLstd450SMax; return GLSLstd450SMax;
} }
case semantic::Intrinsic::kMin: case IntrinsicType::kMin:
if (type->is_float_scalar_or_vector()) { if (type->is_float_scalar_or_vector()) {
return GLSLstd450NMin; return GLSLstd450NMin;
} else if (type->is_unsigned_scalar_or_vector()) { } else if (type->is_unsigned_scalar_or_vector()) {
@ -243,45 +244,45 @@ uint32_t intrinsic_to_glsl_method(type::Type* type,
} else { } else {
return GLSLstd450SMin; return GLSLstd450SMin;
} }
case semantic::Intrinsic::kMix: case IntrinsicType::kMix:
return GLSLstd450FMix; return GLSLstd450FMix;
case semantic::Intrinsic::kModf: case IntrinsicType::kModf:
return GLSLstd450Modf; return GLSLstd450Modf;
case semantic::Intrinsic::kNormalize: case IntrinsicType::kNormalize:
return GLSLstd450Normalize; return GLSLstd450Normalize;
case semantic::Intrinsic::kPack4x8Snorm: case IntrinsicType::kPack4x8Snorm:
return GLSLstd450PackSnorm4x8; return GLSLstd450PackSnorm4x8;
case semantic::Intrinsic::kPack4x8Unorm: case IntrinsicType::kPack4x8Unorm:
return GLSLstd450PackUnorm4x8; return GLSLstd450PackUnorm4x8;
case semantic::Intrinsic::kPack2x16Snorm: case IntrinsicType::kPack2x16Snorm:
return GLSLstd450PackSnorm2x16; return GLSLstd450PackSnorm2x16;
case semantic::Intrinsic::kPack2x16Unorm: case IntrinsicType::kPack2x16Unorm:
return GLSLstd450PackUnorm2x16; return GLSLstd450PackUnorm2x16;
case semantic::Intrinsic::kPack2x16Float: case IntrinsicType::kPack2x16Float:
return GLSLstd450PackHalf2x16; return GLSLstd450PackHalf2x16;
case semantic::Intrinsic::kPow: case IntrinsicType::kPow:
return GLSLstd450Pow; return GLSLstd450Pow;
case semantic::Intrinsic::kReflect: case IntrinsicType::kReflect:
return GLSLstd450Reflect; return GLSLstd450Reflect;
case semantic::Intrinsic::kRound: case IntrinsicType::kRound:
return GLSLstd450Round; return GLSLstd450Round;
case semantic::Intrinsic::kSign: case IntrinsicType::kSign:
return GLSLstd450FSign; return GLSLstd450FSign;
case semantic::Intrinsic::kSin: case IntrinsicType::kSin:
return GLSLstd450Sin; return GLSLstd450Sin;
case semantic::Intrinsic::kSinh: case IntrinsicType::kSinh:
return GLSLstd450Sinh; return GLSLstd450Sinh;
case semantic::Intrinsic::kSmoothStep: case IntrinsicType::kSmoothStep:
return GLSLstd450SmoothStep; return GLSLstd450SmoothStep;
case semantic::Intrinsic::kSqrt: case IntrinsicType::kSqrt:
return GLSLstd450Sqrt; return GLSLstd450Sqrt;
case semantic::Intrinsic::kStep: case IntrinsicType::kStep:
return GLSLstd450Step; return GLSLstd450Step;
case semantic::Intrinsic::kTan: case IntrinsicType::kTan:
return GLSLstd450Tan; return GLSLstd450Tan;
case semantic::Intrinsic::kTanh: case IntrinsicType::kTanh:
return GLSLstd450Tanh; return GLSLstd450Tanh;
case semantic::Intrinsic::kTrunc: case IntrinsicType::kTrunc:
return GLSLstd450Trunc; return GLSLstd450Trunc;
default: default:
break; break;
@ -1897,11 +1898,11 @@ uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident,
OperandList params = {Operand::Int(result_type_id), result}; OperandList params = {Operand::Int(result_type_id), result};
spv::Op op = spv::Op::OpNop; spv::Op op = spv::Op::OpNop;
if (intrinsic == semantic::Intrinsic::kAny) { if (intrinsic == IntrinsicType::kAny) {
op = spv::Op::OpAny; op = spv::Op::OpAny;
} else if (intrinsic == semantic::Intrinsic::kAll) { } else if (intrinsic == IntrinsicType::kAll) {
op = spv::Op::OpAll; op = spv::Op::OpAll;
} else if (intrinsic == semantic::Intrinsic::kArrayLength) { } else if (intrinsic == IntrinsicType::kArrayLength) {
if (call->params().empty()) { if (call->params().empty()) {
error_ = "missing param for runtime array length"; error_ = "missing param for runtime array length";
return 0; return 0;
@ -1934,35 +1935,35 @@ uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident,
return 0; return 0;
} }
return result_id; return result_id;
} else if (intrinsic == semantic::Intrinsic::kCountOneBits) { } else if (intrinsic == IntrinsicType::kCountOneBits) {
op = spv::Op::OpBitCount; op = spv::Op::OpBitCount;
} else if (intrinsic == semantic::Intrinsic::kDot) { } else if (intrinsic == IntrinsicType::kDot) {
op = spv::Op::OpDot; op = spv::Op::OpDot;
} else if (intrinsic == semantic::Intrinsic::kDpdx) { } else if (intrinsic == IntrinsicType::kDpdx) {
op = spv::Op::OpDPdx; op = spv::Op::OpDPdx;
} else if (intrinsic == semantic::Intrinsic::kDpdxCoarse) { } else if (intrinsic == IntrinsicType::kDpdxCoarse) {
op = spv::Op::OpDPdxCoarse; op = spv::Op::OpDPdxCoarse;
} else if (intrinsic == semantic::Intrinsic::kDpdxFine) { } else if (intrinsic == IntrinsicType::kDpdxFine) {
op = spv::Op::OpDPdxFine; op = spv::Op::OpDPdxFine;
} else if (intrinsic == semantic::Intrinsic::kDpdy) { } else if (intrinsic == IntrinsicType::kDpdy) {
op = spv::Op::OpDPdy; op = spv::Op::OpDPdy;
} else if (intrinsic == semantic::Intrinsic::kDpdyCoarse) { } else if (intrinsic == IntrinsicType::kDpdyCoarse) {
op = spv::Op::OpDPdyCoarse; op = spv::Op::OpDPdyCoarse;
} else if (intrinsic == semantic::Intrinsic::kDpdyFine) { } else if (intrinsic == IntrinsicType::kDpdyFine) {
op = spv::Op::OpDPdyFine; op = spv::Op::OpDPdyFine;
} else if (intrinsic == semantic::Intrinsic::kFwidth) { } else if (intrinsic == IntrinsicType::kFwidth) {
op = spv::Op::OpFwidth; op = spv::Op::OpFwidth;
} else if (intrinsic == semantic::Intrinsic::kFwidthCoarse) { } else if (intrinsic == IntrinsicType::kFwidthCoarse) {
op = spv::Op::OpFwidthCoarse; op = spv::Op::OpFwidthCoarse;
} else if (intrinsic == semantic::Intrinsic::kFwidthFine) { } else if (intrinsic == IntrinsicType::kFwidthFine) {
op = spv::Op::OpFwidthFine; op = spv::Op::OpFwidthFine;
} else if (intrinsic == semantic::Intrinsic::kIsInf) { } else if (intrinsic == IntrinsicType::kIsInf) {
op = spv::Op::OpIsInf; op = spv::Op::OpIsInf;
} else if (intrinsic == semantic::Intrinsic::kIsNan) { } else if (intrinsic == IntrinsicType::kIsNan) {
op = spv::Op::OpIsNan; op = spv::Op::OpIsNan;
} else if (intrinsic == semantic::Intrinsic::kReverseBits) { } else if (intrinsic == IntrinsicType::kReverseBits) {
op = spv::Op::OpBitReverse; op = spv::Op::OpBitReverse;
} else if (intrinsic == semantic::Intrinsic::kSelect) { } else if (intrinsic == IntrinsicType::kSelect) {
op = spv::Op::OpSelect; op = spv::Op::OpSelect;
} else { } else {
GenerateGLSLstd450Import(); GenerateGLSLstd450Import();
@ -2163,7 +2164,7 @@ bool Builder::GenerateTextureIntrinsic(
}; };
switch (sem->intrinsic()) { switch (sem->intrinsic()) {
case semantic::Intrinsic::kTextureDimensions: { case IntrinsicType::kTextureDimensions: {
// Number of returned elements from OpImageQuerySize[Lod] may not match // Number of returned elements from OpImageQuerySize[Lod] may not match
// those of textureDimensions(). // those of textureDimensions().
// This might be due to an extra vector scalar describing the number of // This might be due to an extra vector scalar describing the number of
@ -2219,7 +2220,7 @@ bool Builder::GenerateTextureIntrinsic(
} }
break; break;
} }
case semantic::Intrinsic::kTextureNumLayers: { case IntrinsicType::kTextureNumLayers: {
uint32_t spirv_dims = 0; uint32_t spirv_dims = 0;
switch (texture_type->dim()) { switch (texture_type->dim()) {
default: default:
@ -2254,19 +2255,19 @@ bool Builder::GenerateTextureIntrinsic(
} }
break; break;
} }
case semantic::Intrinsic::kTextureNumLevels: { case IntrinsicType::kTextureNumLevels: {
op = spv::Op::OpImageQueryLevels; op = spv::Op::OpImageQueryLevels;
append_result_type_and_id_to_spirv_params(); append_result_type_and_id_to_spirv_params();
spirv_params.emplace_back(gen_param(pidx.texture)); spirv_params.emplace_back(gen_param(pidx.texture));
break; break;
} }
case semantic::Intrinsic::kTextureNumSamples: { case IntrinsicType::kTextureNumSamples: {
op = spv::Op::OpImageQuerySamples; op = spv::Op::OpImageQuerySamples;
append_result_type_and_id_to_spirv_params(); append_result_type_and_id_to_spirv_params();
spirv_params.emplace_back(gen_param(pidx.texture)); spirv_params.emplace_back(gen_param(pidx.texture));
break; break;
} }
case semantic::Intrinsic::kTextureLoad: { case IntrinsicType::kTextureLoad: {
op = texture_type->Is<type::StorageTexture>() ? spv::Op::OpImageRead op = texture_type->Is<type::StorageTexture>() ? spv::Op::OpImageRead
: spv::Op::OpImageFetch; : spv::Op::OpImageFetch;
append_result_type_and_id_to_spirv_params_for_read(); append_result_type_and_id_to_spirv_params_for_read();
@ -2287,7 +2288,7 @@ bool Builder::GenerateTextureIntrinsic(
break; break;
} }
case semantic::Intrinsic::kTextureStore: { case IntrinsicType::kTextureStore: {
op = spv::Op::OpImageWrite; op = spv::Op::OpImageWrite;
spirv_params.emplace_back(gen_param(pidx.texture)); spirv_params.emplace_back(gen_param(pidx.texture));
if (!append_coords_to_spirv_params()) { if (!append_coords_to_spirv_params()) {
@ -2296,7 +2297,7 @@ bool Builder::GenerateTextureIntrinsic(
spirv_params.emplace_back(gen_param(pidx.value)); spirv_params.emplace_back(gen_param(pidx.value));
break; break;
} }
case semantic::Intrinsic::kTextureSample: { case IntrinsicType::kTextureSample: {
op = spv::Op::OpImageSampleImplicitLod; op = spv::Op::OpImageSampleImplicitLod;
append_result_type_and_id_to_spirv_params_for_read(); append_result_type_and_id_to_spirv_params_for_read();
if (!append_image_and_coords_to_spirv_params()) { if (!append_image_and_coords_to_spirv_params()) {
@ -2304,7 +2305,7 @@ bool Builder::GenerateTextureIntrinsic(
} }
break; break;
} }
case semantic::Intrinsic::kTextureSampleBias: { case IntrinsicType::kTextureSampleBias: {
op = spv::Op::OpImageSampleImplicitLod; op = spv::Op::OpImageSampleImplicitLod;
append_result_type_and_id_to_spirv_params_for_read(); append_result_type_and_id_to_spirv_params_for_read();
if (!append_image_and_coords_to_spirv_params()) { if (!append_image_and_coords_to_spirv_params()) {
@ -2315,7 +2316,7 @@ bool Builder::GenerateTextureIntrinsic(
ImageOperand{SpvImageOperandsBiasMask, gen_param(pidx.bias)}); ImageOperand{SpvImageOperandsBiasMask, gen_param(pidx.bias)});
break; break;
} }
case semantic::Intrinsic::kTextureSampleLevel: { case IntrinsicType::kTextureSampleLevel: {
op = spv::Op::OpImageSampleExplicitLod; op = spv::Op::OpImageSampleExplicitLod;
append_result_type_and_id_to_spirv_params_for_read(); append_result_type_and_id_to_spirv_params_for_read();
if (!append_image_and_coords_to_spirv_params()) { if (!append_image_and_coords_to_spirv_params()) {
@ -2339,7 +2340,7 @@ bool Builder::GenerateTextureIntrinsic(
image_operands.emplace_back(ImageOperand{SpvImageOperandsLodMask, level}); image_operands.emplace_back(ImageOperand{SpvImageOperandsLodMask, level});
break; break;
} }
case semantic::Intrinsic::kTextureSampleGrad: { case IntrinsicType::kTextureSampleGrad: {
op = spv::Op::OpImageSampleExplicitLod; op = spv::Op::OpImageSampleExplicitLod;
append_result_type_and_id_to_spirv_params_for_read(); append_result_type_and_id_to_spirv_params_for_read();
if (!append_image_and_coords_to_spirv_params()) { if (!append_image_and_coords_to_spirv_params()) {
@ -2353,7 +2354,7 @@ bool Builder::GenerateTextureIntrinsic(
ImageOperand{SpvImageOperandsGradMask, gen_param(pidx.ddy)}); ImageOperand{SpvImageOperandsGradMask, gen_param(pidx.ddy)});
break; break;
} }
case semantic::Intrinsic::kTextureSampleCompare: { case IntrinsicType::kTextureSampleCompare: {
op = spv::Op::OpImageSampleDrefExplicitLod; op = spv::Op::OpImageSampleDrefExplicitLod;
append_result_type_and_id_to_spirv_params(); append_result_type_and_id_to_spirv_params();
if (!append_image_and_coords_to_spirv_params()) { if (!append_image_and_coords_to_spirv_params()) {