diff --git a/src/semantic/call.h b/src/semantic/call.h index 6351421139..0ef874f785 100644 --- a/src/semantic/call.h +++ b/src/semantic/call.h @@ -26,96 +26,17 @@ namespace semantic { class Call : public Castable { public: /// Constructor - /// @param return_type the return type of the call - explicit Call(type::Type* return_type); + /// @param target the call target + explicit Call(const CallTarget* target); /// Destructor ~Call() override; -}; -/// IntrinsicCall holds semantic information for ast::CallExpression nodes that -/// call intrinsic functions. -class IntrinsicCall : public Castable { - public: - /// Constructor - /// @param return_type the return type of the call - /// @param intrinsic the call target intrinsic - IntrinsicCall(type::Type* return_type, IntrinsicType intrinsic); - - /// Destructor - ~IntrinsicCall() override; - - /// @returns the target intrinsic for the call - IntrinsicType intrinsic() const { return intrinsic_; } + /// @return the target of the call + const CallTarget* Target() const { return target_; } private: - IntrinsicType const intrinsic_; -}; - -/// TextureIntrinsicCall holds semantic information for ast::CallExpression -/// nodes that call intrinsic texture functions. -class TextureIntrinsicCall - : public Castable { - public: - /// Parameters describes the parameters for the texture function. - struct Parameters { - /// kNotUsed is the constant that indicates the given parameter is not part - /// of the texture function signature. - static constexpr const size_t kNotUsed = ~static_cast(0u); - /// Index holds each of the possible parameter indices. If a parameter index - /// is equal to `kNotUsed` then this parameter is not used by the function. - struct Index { - /// Constructor - Index(); - /// Copy constructor - Index(const Index&); - /// `array_index` parameter index. - size_t array_index = kNotUsed; - /// `bias` parameter index. - size_t bias = kNotUsed; - /// `coords` parameter index. - size_t coords = kNotUsed; - /// `depth_ref` parameter index. - size_t depth_ref = kNotUsed; - /// `ddx` parameter index. - size_t ddx = kNotUsed; - /// `ddy` parameter index. - size_t ddy = kNotUsed; - /// `level` parameter index. - size_t level = kNotUsed; - /// `offset` parameter index. - size_t offset = kNotUsed; - /// `sampler` parameter index. - size_t sampler = kNotUsed; - /// `sample_index` parameter index. - size_t sample_index = kNotUsed; - /// `texture` parameter index. - size_t texture = kNotUsed; - /// `value` parameter index. - size_t value = kNotUsed; - }; - /// The indices of all possible parameters. - Index idx; - /// Total number of parameters. - size_t count = 0; - }; - - /// Constructor - /// @param return_type the return type of the call - /// @param intrinsic the call target intrinsic - /// @param params the overload parameter info - TextureIntrinsicCall(type::Type* return_type, - IntrinsicType intrinsic, - const Parameters& params); - - /// Destructor - ~TextureIntrinsicCall() override; - - /// @return the texture call's parameters - const Parameters& Params() const { return params_; } - - private: - const Parameters params_; + CallTarget const* const target_; }; } // namespace semantic diff --git a/src/semantic/call_target.h b/src/semantic/call_target.h index ddb6aacb4f..43cbc0e469 100644 --- a/src/semantic/call_target.h +++ b/src/semantic/call_target.h @@ -32,18 +32,52 @@ namespace semantic { /// Parameter describes a single parameter of a call target struct Parameter { + /// Usage is extra metadata for identifying a parameter based on its overload + /// position + enum class Usage { + kNone, + kArrayIndex, + kBias, + kCoords, + kDepthRef, + kDdx, + kDdy, + kLevel, + kOffset, + kSampler, + kSampleIndex, + kTexture, + kValue, + }; + /// Parameter type - type::Type* type; + type::Type* const type; + /// Parameter usage + Usage const usage = Usage::kNone; }; +/// @returns a string representation of the given parameter usage. +const char* str(Parameter::Usage usage); + +/// Parameters is a list of Parameter using Parameters = std::vector; +/// @param parameters the list of parameters +/// @param usage the parameter usage to find +/// @returns the index of the parameter with the given usage, or -1 if no +/// parameter with the given usage exists. +int IndexOf(const Parameters& parameters, Parameter::Usage usage); + /// CallTarget is the base for callable functions class CallTarget : public Castable { public: /// Constructor + /// @param return_type the return type of the call target /// @param parameters the parameters for the call target - explicit CallTarget(const semantic::Parameters& parameters); + CallTarget(type::Type* return_type, const semantic::Parameters& parameters); + + /// @return the return type of the call target + type::Type* ReturnType() const { return return_type_; } /// Destructor ~CallTarget() override; @@ -52,7 +86,8 @@ class CallTarget : public Castable { const Parameters& Parameters() const { return parameters_; } private: - semantic::Parameters parameters_; + type::Type* const return_type_; + semantic::Parameters const parameters_; }; } // namespace semantic diff --git a/src/semantic/intrinsic.h b/src/semantic/intrinsic.h index f006f31ef2..e2c1902f30 100644 --- a/src/semantic/intrinsic.h +++ b/src/semantic/intrinsic.h @@ -17,6 +17,8 @@ #include +#include "src/semantic/call_target.h" + namespace tint { namespace semantic { @@ -102,39 +104,37 @@ enum class IntrinsicType { kTrunc }; -/// Emits the name of the intrinsic function. The spelling, -/// including case, matches the name in the WGSL spec. -std::ostream& operator<<(std::ostream& out, IntrinsicType i); - -namespace intrinsic { +/// @returns the name of the intrinsic function type. The spelling, including +/// case, matches the name in the WGSL spec. +const char* str(IntrinsicType i); /// Determines if the given `i` is a coarse derivative -/// @param i the intrinsic +/// @param i the intrinsic type /// @returns true if the given derivative is coarse. -bool IsCoarseDerivative(IntrinsicType i); +bool IsCoarseDerivativeIntrinsic(IntrinsicType i); /// Determines if the given `i` is a fine derivative -/// @param i the intrinsic +/// @param i the intrinsic type /// @returns true if the given derivative is fine. -bool IsFineDerivative(IntrinsicType i); +bool IsFineDerivativeIntrinsic(IntrinsicType i); /// Determine if the given `i` is a derivative intrinsic -/// @param i the intrinsic +/// @param i the intrinsic type /// @returns true if the given `i` is a derivative intrinsic -bool IsDerivative(IntrinsicType i); +bool IsDerivativeIntrinsic(IntrinsicType i); /// Determines if the given `i` is a float classification intrinsic -/// @param i the intrinsic +/// @param i the intrinsic type /// @returns true if the given `i` is a float intrinsic bool IsFloatClassificationIntrinsic(IntrinsicType i); /// Determines if the given `i` is a texture operation intrinsic -/// @param i the intrinsic +/// @param i the intrinsic type /// @returns true if the given `i` is a texture operation intrinsic bool IsTextureIntrinsic(IntrinsicType i); /// Determines if the given `i` is a image query intrinsic -/// @param i the intrinsic +/// @param i the intrinsic type /// @returns true if the given `i` is a image query intrinsic bool IsImageQueryIntrinsic(IntrinsicType i); @@ -143,11 +143,56 @@ bool IsImageQueryIntrinsic(IntrinsicType i); /// @returns true if the given `i` is a data packing intrinsic bool IsDataPackingIntrinsic(IntrinsicType i); -/// @returns the name of the intrinsic function. The spelling, including case, -/// matches the name in the WGSL spec. -const char* str(IntrinsicType i); +/// Intrinsic holds the semantic information for an intrinsic function. +class Intrinsic : public Castable { + public: + /// Constructor + /// @param type the intrinsic type + /// @param return_type the return type for the intrinsic call + /// @param parameters the parameters for the intrinsic overload + Intrinsic(IntrinsicType type, + type::Type* return_type, + const semantic::Parameters& parameters); + + /// Destructor + ~Intrinsic() override; + + /// @return the type of the intrinsic + IntrinsicType Type() const { return type_; } + + /// @returns the name of the intrinsic function type. The spelling, including + /// case, matches the name in the WGSL spec. + const char* str() const; + + /// @returns true if intrinsic is a coarse derivative intrinsic + bool IsCoarseDerivative() const; + + /// @returns true if intrinsic is a fine a derivative intrinsic + bool IsFineDerivative() const; + + /// @returns true if intrinsic is a derivative intrinsic + bool IsDerivative() const; + + /// @returns true if intrinsic is a float intrinsic + bool IsFloatClassification() const; + + /// @returns true if intrinsic is a texture operation intrinsic + bool IsTexture() const; + + /// @returns true if intrinsic is a image query intrinsic + bool IsImageQuery() const; + + /// @returns true if intrinsic is a data packing intrinsic + bool IsDataPacking() const; + + private: + IntrinsicType const type_; +}; + +/// Emits the name of the intrinsic function type. The spelling, including case, +/// matches the name in the WGSL spec. +std::ostream& operator<<(std::ostream& out, IntrinsicType i); -} // namespace intrinsic } // namespace semantic } // namespace tint diff --git a/src/semantic/sem_call.cc b/src/semantic/sem_call.cc index 777ad51630..632f3a0656 100644 --- a/src/semantic/sem_call.cc +++ b/src/semantic/sem_call.cc @@ -15,30 +15,14 @@ #include "src/semantic/call.h" TINT_INSTANTIATE_CLASS_ID(tint::semantic::Call); -TINT_INSTANTIATE_CLASS_ID(tint::semantic::IntrinsicCall); -TINT_INSTANTIATE_CLASS_ID(tint::semantic::TextureIntrinsicCall); namespace tint { namespace semantic { -Call::Call(type::Type* return_type) : Base(return_type) {} +Call::Call(const CallTarget* target) + : Base(target->ReturnType()), target_(target) {} Call::~Call() = default; -IntrinsicCall::IntrinsicCall(type::Type* return_type, IntrinsicType intrinsic) - : Base(return_type), intrinsic_(intrinsic) {} - -IntrinsicCall::~IntrinsicCall() = default; - -TextureIntrinsicCall::TextureIntrinsicCall(type::Type* return_type, - IntrinsicType intrinsic, - const Parameters& params) - : Base(return_type, intrinsic), params_(params) {} - -TextureIntrinsicCall::~TextureIntrinsicCall() = default; - -TextureIntrinsicCall::Parameters::Index::Index() = default; -TextureIntrinsicCall::Parameters::Index::Index(const Index&) = default; - } // namespace semantic } // namespace tint diff --git a/src/semantic/sem_call_target.cc b/src/semantic/sem_call_target.cc index 3b79a555cb..ad7f30070a 100644 --- a/src/semantic/sem_call_target.cc +++ b/src/semantic/sem_call_target.cc @@ -21,10 +21,50 @@ TINT_INSTANTIATE_CLASS_ID(tint::semantic::CallTarget); namespace tint { namespace semantic { -CallTarget::CallTarget(const semantic::Parameters& parameters) - : parameters_(parameters) {} +CallTarget::CallTarget(type::Type* return_type, + const semantic::Parameters& parameters) + : return_type_(return_type), parameters_(parameters) {} CallTarget::~CallTarget() = default; +int IndexOf(const Parameters& parameters, Parameter::Usage usage) { + for (size_t i = 0; i < parameters.size(); i++) { + if (parameters[i].usage == usage) { + return static_cast(i); + } + } + return -1; +} + +const char* str(Parameter::Usage usage) { + switch (usage) { + case Parameter::Usage::kArrayIndex: + return "array_index"; + case Parameter::Usage::kBias: + return "bias"; + case Parameter::Usage::kCoords: + return "coords"; + case Parameter::Usage::kDepthRef: + return "depth_ref"; + case Parameter::Usage::kDdx: + return "ddx"; + case Parameter::Usage::kDdy: + return "ddy"; + case Parameter::Usage::kLevel: + return "level"; + case Parameter::Usage::kOffset: + return "offset"; + case Parameter::Usage::kSampler: + return "sampler"; + case Parameter::Usage::kSampleIndex: + return "sample_index"; + case Parameter::Usage::kTexture: + return "texture"; + case Parameter::Usage::kValue: + return "value"; + default: + return ""; + } +} } // namespace semantic } // namespace tint diff --git a/src/semantic/sem_function.cc b/src/semantic/sem_function.cc index 8c8d16c0eb..db533398eb 100644 --- a/src/semantic/sem_function.cc +++ b/src/semantic/sem_function.cc @@ -37,7 +37,7 @@ Parameters GetParameters(ast::Function* ast) { semantic::Parameters parameters; parameters.reserve(ast->params().size()); for (auto* param : ast->params()) { - parameters.emplace_back(Parameter{param->type()}); + parameters.emplace_back(Parameter{param->type(), Parameter::Usage::kNone}); } return parameters; } @@ -48,7 +48,7 @@ Function::Function(ast::Function* ast, std::vector referenced_module_vars, std::vector local_referenced_module_vars, std::vector ancestor_entry_points) - : Base(GetParameters(ast)), + : Base(ast->return_type(), GetParameters(ast)), referenced_module_vars_(std::move(referenced_module_vars)), local_referenced_module_vars_(std::move(local_referenced_module_vars)), ancestor_entry_points_(std::move(ancestor_entry_points)) {} diff --git a/src/semantic/sem_intrinsic.cc b/src/semantic/sem_intrinsic.cc index b2eea277e7..e12c1b4130 100644 --- a/src/semantic/sem_intrinsic.cc +++ b/src/semantic/sem_intrinsic.cc @@ -14,15 +14,19 @@ #include "src/semantic/intrinsic.h" +TINT_INSTANTIATE_CLASS_ID(tint::semantic::Intrinsic); + namespace tint { namespace semantic { std::ostream& operator<<(std::ostream& out, IntrinsicType i) { - out << intrinsic::str(i); + out << str(i); return out; } -namespace intrinsic { +const char* Intrinsic::str() const { + return semantic::str(type_); +} const char* str(IntrinsicType i) { /// The emitted name matches the spelling in the WGSL spec. @@ -188,20 +192,20 @@ const char* str(IntrinsicType i) { return ""; } -bool IsCoarseDerivative(IntrinsicType i) { +bool IsCoarseDerivativeIntrinsic(IntrinsicType i) { return i == IntrinsicType::kDpdxCoarse || i == IntrinsicType::kDpdyCoarse || i == IntrinsicType::kFwidthCoarse; } -bool IsFineDerivative(IntrinsicType i) { +bool IsFineDerivativeIntrinsic(IntrinsicType i) { return i == IntrinsicType::kDpdxFine || i == IntrinsicType::kDpdyFine || i == IntrinsicType::kFwidthFine; } -bool IsDerivative(IntrinsicType i) { +bool IsDerivativeIntrinsic(IntrinsicType i) { return i == IntrinsicType::kDpdx || i == IntrinsicType::kDpdy || - i == IntrinsicType::kFwidth || IsCoarseDerivative(i) || - IsFineDerivative(i); + i == IntrinsicType::kFwidth || IsCoarseDerivativeIntrinsic(i) || + IsFineDerivativeIntrinsic(i); } bool IsFloatClassificationIntrinsic(IntrinsicType i) { @@ -220,7 +224,7 @@ bool IsTextureIntrinsic(IntrinsicType i) { } bool IsImageQueryIntrinsic(IntrinsicType i) { - return i == semantic::IntrinsicType::kTextureDimensions || + return i == IntrinsicType::kTextureDimensions || i == IntrinsicType::kTextureNumLayers || i == IntrinsicType::kTextureNumLevels || i == IntrinsicType::kTextureNumSamples; @@ -234,6 +238,40 @@ bool IsDataPackingIntrinsic(IntrinsicType i) { i == IntrinsicType::kPack2x16Float; } -} // namespace intrinsic +Intrinsic::Intrinsic(IntrinsicType type, + type::Type* return_type, + const semantic::Parameters& parameters) + : Base(return_type, parameters), type_(type) {} + +Intrinsic::~Intrinsic() = default; + +bool Intrinsic::IsCoarseDerivative() const { + return IsCoarseDerivativeIntrinsic(type_); +} + +bool Intrinsic::IsFineDerivative() const { + return IsFineDerivativeIntrinsic(type_); +} + +bool Intrinsic::IsDerivative() const { + return IsDerivativeIntrinsic(type_); +} + +bool Intrinsic::IsFloatClassification() const { + return IsFloatClassificationIntrinsic(type_); +} + +bool Intrinsic::IsTexture() const { + return IsTextureIntrinsic(type_); +} + +bool Intrinsic::IsImageQuery() const { + return IsImageQueryIntrinsic(type_); +} + +bool Intrinsic::IsDataPacking() const { + return IsDataPackingIntrinsic(type_); +} + } // namespace semantic } // namespace tint diff --git a/src/type_determiner.cc b/src/type_determiner.cc index 6b47ffa587..7ea3c093fd 100644 --- a/src/type_determiner.cc +++ b/src/type_determiner.cc @@ -419,9 +419,9 @@ bool TypeDeterminer::DetermineCall(ast::CallExpression* call) { auto name = builder_->Symbols().NameFor(ident->symbol()); - auto intrinsic = MatchIntrinsic(name); - if (intrinsic != IntrinsicType::kNone) { - if (!DetermineIntrinsicCall(call, intrinsic)) { + auto intrinsic_type = MatchIntrinsicType(name); + if (intrinsic_type != IntrinsicType::kNone) { + if (!DetermineIntrinsicCall(call, intrinsic_type)) { return false; } } else { @@ -450,9 +450,7 @@ bool TypeDeterminer::DetermineCall(ast::CallExpression* call) { } auto* function = iter->second; - auto* return_ty = function->declaration->return_type(); - auto* sem = builder_->create(return_ty); - builder_->Sem().Add(call, sem); + function_calls_.emplace(call, function); } return true; @@ -548,14 +546,26 @@ constexpr const uint32_t kIntrinsicDataCount = } // namespace bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call, - IntrinsicType intrinsic) { - auto create_sem = [&](type::Type* result) { - auto* sem = builder_->create(result, intrinsic); - builder_->Sem().Add(call, sem); + IntrinsicType intrinsic_type) { + using Parameter = semantic::Parameter; + using Parameters = semantic::Parameters; + using Usage = Parameter::Usage; + + std::vector arg_tys; + arg_tys.reserve(call->params().size()); + for (auto* expr : call->params()) { + arg_tys.emplace_back(TypeOf(expr)); + } + + auto create_sem = [&](type::Type* return_type) { + semantic::Parameters params; // TODO(bclayton): Populate this + auto* intrinsic = builder_->create( + intrinsic_type, return_type, params); + builder_->Sem().Add(call, builder_->create(intrinsic)); }; - std::string name = semantic::intrinsic::str(intrinsic); - if (semantic::intrinsic::IsFloatClassificationIntrinsic(intrinsic)) { + std::string name = semantic::str(intrinsic_type); + if (semantic::IsFloatClassificationIntrinsic(intrinsic_type)) { if (call->params().size() != 1) { set_error(call->source(), "incorrect number of parameters for " + name); return false; @@ -571,131 +581,135 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call, } return true; } - if (semantic::intrinsic::IsTextureIntrinsic(intrinsic)) { - semantic::TextureIntrinsicCall::Parameters param; + if (semantic::IsTextureIntrinsic(intrinsic_type)) { + Parameters params; - auto* texture_param = call->params()[0]; - if (!TypeOf(texture_param)->UnwrapAll()->Is()) { + auto& ty = builder_->ty; + + auto* texture = arg_tys[0]->UnwrapAll()->As(); + if (!texture) { set_error(call->source(), "invalid first argument for " + name); return false; } - type::Texture* texture = - TypeOf(texture_param)->UnwrapAll()->As(); bool is_array = type::IsTextureArray(texture->dim()); bool is_multisampled = texture->Is(); - switch (intrinsic) { + switch (intrinsic_type) { case IntrinsicType::kTextureDimensions: - param.idx.texture = param.count++; - if (call->params().size() > param.count) { - param.idx.level = param.count++; + params.emplace_back(Parameter{texture, Usage::kTexture}); + if (arg_tys.size() > params.size()) { + params.emplace_back(Parameter{ty.i32(), Usage::kLevel}); } break; case IntrinsicType::kTextureNumLayers: case IntrinsicType::kTextureNumLevels: case IntrinsicType::kTextureNumSamples: - param.idx.texture = param.count++; + params.emplace_back(Parameter{texture, Usage::kTexture}); break; case IntrinsicType::kTextureLoad: - param.idx.texture = param.count++; - param.idx.coords = param.count++; + params.emplace_back(Parameter{texture, Usage::kTexture}); + params.emplace_back(Parameter{arg_tys[1], Usage::kCoords}); if (is_array) { - param.idx.array_index = param.count++; + params.emplace_back(Parameter{ty.i32(), Usage::kArrayIndex}); } - if (call->params().size() > param.count) { + if (arg_tys.size() > params.size()) { if (is_multisampled) { - param.idx.sample_index = param.count++; + params.emplace_back(Parameter{ty.i32(), Usage::kSampleIndex}); } else { - param.idx.level = param.count++; + params.emplace_back(Parameter{ty.i32(), Usage::kLevel}); } } break; case IntrinsicType::kTextureSample: - param.idx.texture = param.count++; - param.idx.sampler = param.count++; - param.idx.coords = param.count++; + params.emplace_back(Parameter{texture, Usage::kTexture}); + params.emplace_back(Parameter{arg_tys[1], Usage::kSampler}); + params.emplace_back(Parameter{arg_tys[2], Usage::kCoords}); if (is_array) { - param.idx.array_index = param.count++; + params.emplace_back(Parameter{ty.i32(), Usage::kArrayIndex}); } - if (call->params().size() > param.count) { - param.idx.offset = param.count++; + if (arg_tys.size() > params.size()) { + params.emplace_back( + Parameter{arg_tys[params.size()], Usage::kOffset}); } break; case IntrinsicType::kTextureSampleBias: - param.idx.texture = param.count++; - param.idx.sampler = param.count++; - param.idx.coords = param.count++; + params.emplace_back(Parameter{texture, Usage::kTexture}); + params.emplace_back(Parameter{arg_tys[1], Usage::kSampler}); + params.emplace_back(Parameter{arg_tys[2], Usage::kCoords}); if (is_array) { - param.idx.array_index = param.count++; + params.emplace_back(Parameter{ty.i32(), Usage::kArrayIndex}); } - param.idx.bias = param.count++; - if (call->params().size() > param.count) { - param.idx.offset = param.count++; + params.emplace_back(Parameter{ty.f32(), Usage::kBias}); + if (arg_tys.size() > params.size()) { + params.emplace_back( + Parameter{arg_tys[params.size()], Usage::kOffset}); } break; case IntrinsicType::kTextureSampleLevel: - param.idx.texture = param.count++; - param.idx.sampler = param.count++; - param.idx.coords = param.count++; + params.emplace_back(Parameter{texture, Usage::kTexture}); + params.emplace_back(Parameter{arg_tys[1], Usage::kSampler}); + params.emplace_back(Parameter{arg_tys[2], Usage::kCoords}); if (is_array) { - param.idx.array_index = param.count++; + params.emplace_back(Parameter{ty.i32(), Usage::kArrayIndex}); } - param.idx.level = param.count++; - if (call->params().size() > param.count) { - param.idx.offset = param.count++; + params.emplace_back(Parameter{ty.i32(), Usage::kLevel}); + if (arg_tys.size() > params.size()) { + params.emplace_back( + Parameter{arg_tys[params.size()], Usage::kOffset}); } break; case IntrinsicType::kTextureSampleCompare: - param.idx.texture = param.count++; - param.idx.sampler = param.count++; - param.idx.coords = param.count++; + params.emplace_back(Parameter{texture, Usage::kTexture}); + params.emplace_back(Parameter{arg_tys[1], Usage::kSampler}); + params.emplace_back(Parameter{arg_tys[2], Usage::kCoords}); if (is_array) { - param.idx.array_index = param.count++; + params.emplace_back(Parameter{ty.i32(), Usage::kArrayIndex}); } - param.idx.depth_ref = param.count++; - if (call->params().size() > param.count) { - param.idx.offset = param.count++; + params.emplace_back(Parameter{ty.f32(), Usage::kDepthRef}); + if (arg_tys.size() > params.size()) { + params.emplace_back( + Parameter{arg_tys[params.size()], Usage::kOffset}); } break; case IntrinsicType::kTextureSampleGrad: - param.idx.texture = param.count++; - param.idx.sampler = param.count++; - param.idx.coords = param.count++; + params.emplace_back(Parameter{texture, Usage::kTexture}); + params.emplace_back(Parameter{arg_tys[1], Usage::kSampler}); + params.emplace_back(Parameter{arg_tys[2], Usage::kCoords}); if (is_array) { - param.idx.array_index = param.count++; + params.emplace_back(Parameter{ty.i32(), Usage::kArrayIndex}); } - param.idx.ddx = param.count++; - param.idx.ddy = param.count++; - if (call->params().size() > param.count) { - param.idx.offset = param.count++; + params.emplace_back(Parameter{arg_tys[params.size()], Usage::kDdx}); + params.emplace_back(Parameter{arg_tys[params.size()], Usage::kDdy}); + if (arg_tys.size() > params.size()) { + params.emplace_back( + Parameter{arg_tys[params.size()], Usage::kOffset}); } break; case IntrinsicType::kTextureStore: - param.idx.texture = param.count++; - param.idx.coords = param.count++; + params.emplace_back(Parameter{texture, Usage::kTexture}); + params.emplace_back(Parameter{arg_tys[1], Usage::kCoords}); if (is_array) { - param.idx.array_index = param.count++; + params.emplace_back(Parameter{ty.i32(), Usage::kArrayIndex}); } - param.idx.value = param.count++; + params.emplace_back(Parameter{arg_tys[params.size()], Usage::kValue}); break; default: set_error(call->source(), - "Internal compiler error: Unreachable intrinsic " + - std::to_string(static_cast(intrinsic))); + "Internal compiler error: Unreachable intrinsic " + name); return false; } - if (call->params().size() != param.count) { - set_error(call->source(), - "incorrect number of parameters for " + name + ", got " + - std::to_string(call->params().size()) + " and expected " + - std::to_string(param.count)); + if (arg_tys.size() != params.size()) { + set_error(call->source(), "incorrect number of arguments for " + name + + ", got " + std::to_string(arg_tys.size()) + + " and expected " + + std::to_string(params.size())); return false; } // Set the function return type type::Type* return_type = nullptr; - switch (intrinsic) { + switch (intrinsic_type) { case IntrinsicType::kTextureDimensions: { auto* i32 = builder_->create(); switch (texture->dim()) { @@ -748,16 +762,15 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call, } } - auto* sem = builder_->create( - return_type, intrinsic, param); - builder_->Sem().Add(call, sem); - + auto* intrinsic = builder_->create( + intrinsic_type, return_type, params); + builder_->Sem().Add(call, builder_->create(intrinsic)); return true; } const IntrinsicData* data = nullptr; for (uint32_t i = 0; i < kIntrinsicDataCount; ++i) { - if (intrinsic == kIntrinsicData[i].intrinsic) { + if (intrinsic_type == kIntrinsicData[i].intrinsic) { data = &kIntrinsicData[i]; break; } @@ -847,7 +860,7 @@ bool TypeDeterminer::DetermineIdentifier(ast::IdentifierExpression* expr) { } std::string name = builder_->Symbols().NameFor(symbol); - if (MatchIntrinsic(name) != IntrinsicType::kNone) { + if (MatchIntrinsicType(name) != IntrinsicType::kNone) { // Identifier is to an intrinsic function, which has no type (currently). return true; } @@ -857,7 +870,7 @@ bool TypeDeterminer::DetermineIdentifier(ast::IdentifierExpression* expr) { return false; } -IntrinsicType TypeDeterminer::MatchIntrinsic(const std::string& name) { +IntrinsicType TypeDeterminer::MatchIntrinsicType(const std::string& name) { if (name == "abs") { return IntrinsicType::kAbs; } else if (name == "acos") { @@ -1197,14 +1210,23 @@ void TypeDeterminer::CreateSemanticNodes() const { return out; }; + std::unordered_map func_info_to_sem_func; for (auto it : function_to_info_) { auto* func = it.first; auto* info = it.second; - sem.Add(func, - builder_->create( - info->declaration, remap_vars(info->referenced_module_vars), - remap_vars(info->local_referenced_module_vars), - info->ancestor_entry_points)); + auto* sem_func = builder_->create( + info->declaration, remap_vars(info->referenced_module_vars), + remap_vars(info->local_referenced_module_vars), + info->ancestor_entry_points); + func_info_to_sem_func.emplace(info, sem_func); + sem.Add(func, sem_func); + } + + for (auto it : function_calls_) { + auto* call = it.first; + auto* func_info = it.second; + auto* sem_func = func_info_to_sem_func.at(func_info); + builder_->Sem().Add(call, builder_->create(sem_func)); } } diff --git a/src/type_determiner.h b/src/type_determiner.h index c8dc34d297..9928b907ad 100644 --- a/src/type_determiner.h +++ b/src/type_determiner.h @@ -24,6 +24,7 @@ #include "src/diagnostic/diagnostic.h" #include "src/program_builder.h" #include "src/scope_stack.h" +#include "src/semantic/intrinsic.h" #include "src/type/storage_texture_type.h" namespace tint { @@ -66,10 +67,10 @@ class TypeDeterminer { /// @returns true if the type determiner was successful 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 type. /// @return the semantic::IntrinsicType for the given name. If `name` does not - /// match an intrinsic, returns semantic::IntrinsicType::kNone - static semantic::IntrinsicType MatchIntrinsic(const std::string& name); + /// match an intrinsic, returns semantic::Intrinsic::kNone + static semantic::IntrinsicType MatchIntrinsicType(const std::string& name); private: template @@ -177,7 +178,7 @@ class TypeDeterminer { bool DetermineConstructor(ast::ConstructorExpression* expr); bool DetermineIdentifier(ast::IdentifierExpression* expr); bool DetermineIntrinsicCall(ast::CallExpression* call, - semantic::IntrinsicType intrinsic); + semantic::IntrinsicType intrinsic_type); bool DetermineMemberAccessor(ast::MemberAccessorExpression* expr); bool DetermineUnaryOp(ast::UnaryOpExpression* expr); @@ -195,12 +196,13 @@ class TypeDeterminer { /// @param type the resolved type void SetType(ast::Expression* expr, type::Type* type) const; - ProgramBuilder* builder_; + ProgramBuilder* const builder_; std::string error_; ScopeStack variable_stack_; std::unordered_map symbol_to_function_; std::unordered_map function_to_info_; std::unordered_map variable_to_info_; + std::unordered_map function_calls_; FunctionInfo* current_function_ = nullptr; BlockAllocator variable_infos_; BlockAllocator function_infos_; diff --git a/src/type_determiner_test.cc b/src/type_determiner_test.cc index 72f9e5929d..4b43554ca4 100644 --- a/src/type_determiner_test.cc +++ b/src/type_determiner_test.cc @@ -1637,7 +1637,7 @@ using IntrinsicDataTest = TypeDeterminerTestWithParam; TEST_P(IntrinsicDataTest, Lookup) { auto param = GetParam(); - EXPECT_EQ(TypeDeterminer::MatchIntrinsic(param.name), param.intrinsic); + EXPECT_EQ(TypeDeterminer::MatchIntrinsicType(param.name), param.intrinsic); } INSTANTIATE_TEST_SUITE_P( TypeDeterminerTest, @@ -1717,7 +1717,7 @@ INSTANTIATE_TEST_SUITE_P( IntrinsicData{"trunc", IntrinsicType::kTrunc})); TEST_F(TypeDeterminerTest, MatchIntrinsicNoMatch) { - EXPECT_EQ(TypeDeterminer::MatchIntrinsic("not_intrinsic"), + EXPECT_EQ(TypeDeterminer::MatchIntrinsicType("not_intrinsic"), IntrinsicType::kNone); } @@ -2466,42 +2466,15 @@ INSTANTIATE_TEST_SUITE_P( testing::ValuesIn(ast::intrinsic::test::TextureOverloadCase::ValidCases())); std::string to_str(const std::string& function, - const semantic::TextureIntrinsicCall::Parameters& params) { - struct Parameter { - size_t idx; - std::string name; - }; - std::vector list; - auto maybe_add_param = [&list](size_t idx, const char* name) { - if (idx != - semantic::TextureIntrinsicCall::Parameters::Parameters::kNotUsed) { - list.emplace_back(Parameter{idx, name}); - } - }; - maybe_add_param(params.idx.array_index, "array_index"); - maybe_add_param(params.idx.bias, "bias"); - maybe_add_param(params.idx.coords, "coords"); - maybe_add_param(params.idx.depth_ref, "depth_ref"); - maybe_add_param(params.idx.ddx, "ddx"); - maybe_add_param(params.idx.ddy, "ddy"); - maybe_add_param(params.idx.level, "level"); - maybe_add_param(params.idx.offset, "offset"); - maybe_add_param(params.idx.sampler, "sampler"); - maybe_add_param(params.idx.sample_index, "sample_index"); - maybe_add_param(params.idx.texture, "texture"); - maybe_add_param(params.idx.value, "value"); - std::sort( - list.begin(), list.end(), - [](const Parameter& a, const Parameter& b) { return a.idx < b.idx; }); - + const semantic::Parameters& params) { std::stringstream out; out << function << "("; bool first = true; - for (auto& param : list) { + for (auto& param : params) { if (!first) { out << ", "; } - out << param.name; + out << semantic::str(param.usage); first = false; } out << ")"; @@ -2833,12 +2806,12 @@ TEST_P(TypeDeterminerTextureIntrinsicTest, Call) { } } - auto* sem = Sem().Get(call); - ASSERT_NE(sem, nullptr); - auto* intrinsic = sem->As(); - ASSERT_NE(intrinsic, nullptr); + auto* call_sem = Sem().Get(call); + ASSERT_NE(call_sem, nullptr); + auto* target = call_sem->Target(); + ASSERT_NE(target, nullptr); - auto got = ::tint::to_str(param.function, intrinsic->Params()); + auto got = ::tint::to_str(param.function, target->Parameters()); auto* expected = expected_texture_overload(param.overload); EXPECT_EQ(got, expected); } diff --git a/src/validator/validator_impl.cc b/src/validator/validator_impl.cc index ecad8efef7..019423eb73 100644 --- a/src/validator/validator_impl.cc +++ b/src/validator/validator_impl.cc @@ -644,24 +644,23 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) { return false; } - auto* call_sem = program_->Sem().Get(expr); - if (call_sem == nullptr) { + auto* call = program_->Sem().Get(expr); + if (call == nullptr) { add_error(expr->source(), "CallExpression is missing semantic information"); return false; } - if (auto* intrinsic_sem = call_sem->As()) { + if (auto* intrinsic = call->Target()->As()) { const IntrinsicData* data = nullptr; for (uint32_t i = 0; i < kIntrinsicDataCount; ++i) { - if (intrinsic_sem->intrinsic() == kIntrinsicData[i].intrinsic) { + if (intrinsic->Type() == kIntrinsicData[i].intrinsic) { data = &kIntrinsicData[i]; break; } } if (data != nullptr) { - std::string builtin = - semantic::intrinsic::str(intrinsic_sem->intrinsic()); + std::string builtin = intrinsic->str(); if (expr->params().size() != data->param_count) { add_error(expr->source(), "incorrect number of parameters for " + builtin + @@ -832,7 +831,7 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) { } } - if (semantic::intrinsic::IsDataPackingIntrinsic(data->intrinsic)) { + if (semantic::IsDataPackingIntrinsic(data->intrinsic)) { if (!program_->TypeOf(expr)->Is()) { add_error(expr->source(), "incorrect type for " + builtin + diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index 2a4d304af3..9ac33c8dbb 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -542,22 +542,22 @@ bool GeneratorImpl::EmitCall(std::ostream& pre, return 0; } - auto* call_sem = builder_.Sem().Get(expr); - if (auto* sem = call_sem->As()) { - return EmitTextureCall(pre, out, expr, sem); - } - if (auto* sem = call_sem->As()) { + auto* call = builder_.Sem().Get(expr); + if (auto* intrinsic = call->Target()->As()) { + if (intrinsic->IsTexture()) { + return EmitTextureCall(pre, out, expr, intrinsic); + } const auto& params = expr->params(); - if (sem->intrinsic() == semantic::IntrinsicType::kSelect) { + if (intrinsic->Type() == semantic::IntrinsicType::kSelect) { error_ = "select not supported in HLSL backend yet"; return false; - } else if (sem->intrinsic() == semantic::IntrinsicType::kIsNormal) { + } else if (intrinsic->Type() == semantic::IntrinsicType::kIsNormal) { error_ = "is_normal not supported in HLSL backend yet"; return false; - } else if (semantic::intrinsic::IsDataPackingIntrinsic(sem->intrinsic())) { - return EmitDataPackingCall(pre, out, expr); + } else if (intrinsic->IsDataPacking()) { + return EmitDataPackingCall(pre, out, expr, intrinsic); } - auto name = generate_builtin_name(sem); + auto name = generate_builtin_name(intrinsic); if (name.empty()) { return false; } @@ -638,9 +638,8 @@ bool GeneratorImpl::EmitCall(std::ostream& pre, bool GeneratorImpl::EmitDataPackingCall(std::ostream& pre, std::ostream& out, - ast::CallExpression* expr) { - auto* ident = builder_.Sem().Get(expr)->As(); - + ast::CallExpression* expr, + const semantic::Intrinsic* intrinsic) { auto* param = expr->params()[0]; auto tmp_name = generate_name(kTempNamePrefix); std::ostringstream expr_out; @@ -650,17 +649,17 @@ bool GeneratorImpl::EmitDataPackingCall(std::ostream& pre, uint32_t dims = 2; bool is_signed = false; uint32_t scale = 65535; - if (ident->intrinsic() == semantic::IntrinsicType::kPack4x8Snorm || - ident->intrinsic() == semantic::IntrinsicType::kPack4x8Unorm) { + if (intrinsic->Type() == semantic::IntrinsicType::kPack4x8Snorm || + intrinsic->Type() == semantic::IntrinsicType::kPack4x8Unorm) { dims = 4; scale = 255; } - if (ident->intrinsic() == semantic::IntrinsicType::kPack4x8Snorm || - ident->intrinsic() == semantic::IntrinsicType::kPack2x16Snorm) { + if (intrinsic->Type() == semantic::IntrinsicType::kPack4x8Snorm || + intrinsic->Type() == semantic::IntrinsicType::kPack2x16Snorm) { is_signed = true; scale = (scale - 1) / 2; } - switch (ident->intrinsic()) { + switch (intrinsic->Type()) { case semantic::IntrinsicType::kPack4x8Snorm: case semantic::IntrinsicType::kPack4x8Unorm: case semantic::IntrinsicType::kPack2x16Snorm: @@ -698,17 +697,24 @@ bool GeneratorImpl::EmitDataPackingCall(std::ostream& pre, bool GeneratorImpl::EmitTextureCall(std::ostream& pre, std::ostream& out, ast::CallExpression* expr, - const semantic::TextureIntrinsicCall* sem) { - auto* ident = expr->func()->As(); + const semantic::Intrinsic* intrinsic) { + using Usage = semantic::Parameter::Usage; - auto params = expr->params(); - auto& pidx = sem->Params().idx; - auto const kNotUsed = semantic::TextureIntrinsicCall::Parameters::kNotUsed; + auto parameters = intrinsic->Parameters(); + auto arguments = expr->params(); + + // Returns the argument with the given usage + auto arg = [&](Usage usage) { + int idx = semantic::IndexOf(parameters, usage); + return (idx >= 0) ? arguments[idx] : nullptr; + }; + + auto* texture = arg(Usage::kTexture); + assert(texture); - auto* texture = params[pidx.texture]; auto* texture_type = TypeOf(texture)->UnwrapAll()->As(); - switch (sem->intrinsic()) { + switch (intrinsic->Type()) { case semantic::IntrinsicType::kTextureDimensions: case semantic::IntrinsicType::kTextureNumLayers: case semantic::IntrinsicType::kTextureNumLevels: @@ -718,7 +724,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre, const char* swizzle = ""; bool add_mip_level_in = false; - switch (sem->intrinsic()) { + switch (intrinsic->Type()) { case semantic::IntrinsicType::kTextureDimensions: switch (texture_type->dim()) { case type::TextureDimension::kNone: @@ -823,8 +829,11 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre, return false; } pre << ".GetDimensions("; - if (pidx.level != kNotUsed) { - pre << pidx.level << ", "; + if (auto* level = arg(Usage::kLevel)) { + if (!EmitExpression(pre, pre, level)) { + return false; + } + pre << ", "; } else if (add_mip_level_in) { pre << "0, "; } @@ -859,7 +868,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre, bool pack_mip_in_coords = false; - switch (sem->intrinsic()) { + switch (intrinsic->Type()) { case semantic::IntrinsicType::kTextureSample: out << ".Sample("; break; @@ -886,17 +895,18 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre, break; default: error_ = "Internal compiler error: Unhandled texture intrinsic '" + - builder_.Symbols().NameFor(ident->symbol()) + "'"; + std::string(intrinsic->str()) + "'"; return false; } - if (pidx.sampler != kNotUsed) { - if (!EmitExpression(pre, out, params[pidx.sampler])) + if (auto* sampler = arg(Usage::kSampler)) { + if (!EmitExpression(pre, out, sampler)) return false; out << ", "; } - auto* param_coords = params[pidx.coords]; + auto* param_coords = arg(Usage::kCoords); + assert(param_coords); auto emit_vector_appended_with_i32_zero = [&](tint::ast::Expression* vector) { auto* i32 = builder_.create(); @@ -906,10 +916,9 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre, return EmitExpression(pre, out, packed); }; - if (pidx.array_index != kNotUsed) { + if (auto* array_index = arg(Usage::kArrayIndex)) { // Array index needs to be appended to the coordinates. - auto* param_array_index = params[pidx.array_index]; - auto* packed = AppendVector(&builder_, param_coords, param_array_index); + auto* packed = AppendVector(&builder_, param_coords, array_index); if (pack_mip_in_coords) { if (!emit_vector_appended_with_i32_zero(packed)) { return false; @@ -928,18 +937,18 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre, return false; } - for (auto idx : {pidx.depth_ref, pidx.bias, pidx.level, pidx.ddx, pidx.ddy, - pidx.sample_index, pidx.offset}) { - if (idx != kNotUsed) { + for (auto usage : {Usage::kDepthRef, Usage::kBias, Usage::kLevel, Usage::kDdx, + Usage::kDdy, Usage::kSampleIndex, Usage::kOffset}) { + if (auto* e = arg(usage)) { out << ", "; - if (!EmitExpression(pre, out, params[idx])) + if (!EmitExpression(pre, out, e)) return false; } } - if (sem->intrinsic() == semantic::IntrinsicType::kTextureStore) { + if (intrinsic->Type() == semantic::IntrinsicType::kTextureStore) { out << "] = "; - if (!EmitExpression(pre, out, params[pidx.value])) + if (!EmitExpression(pre, out, arg(Usage::kValue))) return false; } else { out << ")"; @@ -949,9 +958,9 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& pre, } // namespace hlsl std::string GeneratorImpl::generate_builtin_name( - const semantic::IntrinsicCall* call) { + const semantic::Intrinsic* intrinsic) { std::string out; - switch (call->intrinsic()) { + switch (intrinsic->Type()) { case semantic::IntrinsicType::kAcos: case semantic::IntrinsicType::kAny: case semantic::IntrinsicType::kAll: @@ -990,7 +999,7 @@ std::string GeneratorImpl::generate_builtin_name( case semantic::IntrinsicType::kMax: case semantic::IntrinsicType::kMin: case semantic::IntrinsicType::kClamp: - out = semantic::intrinsic::str(call->intrinsic()); + out = intrinsic->str(); break; case semantic::IntrinsicType::kCountOneBits: out = "countbits"; @@ -1043,8 +1052,7 @@ std::string GeneratorImpl::generate_builtin_name( out = "smoothstep"; break; default: - error_ = "Unknown builtin method: " + - std::string(semantic::intrinsic::str(call->intrinsic())); + error_ = "Unknown builtin method: " + std::string(intrinsic->str()); return ""; } diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h index 22b8be0545..805a0ceec3 100644 --- a/src/writer/hlsl/generator_impl.h +++ b/src/writer/hlsl/generator_impl.h @@ -48,8 +48,8 @@ namespace tint { // Forward declarations namespace semantic { -class TextureIntrinsicCall; -class IntrinsicCall; +class Call; +class Intrinsic; } // namespace semantic namespace writer { @@ -153,20 +153,22 @@ class GeneratorImpl { /// @param pre the preamble for the expression stream /// @param out the output of the expression stream /// @param expr the call expression - /// @param sem the semantic information for the texture intrinsic call + /// @param intrinsic the semantic information for the texture intrinsic /// @returns true if the call expression is emitted bool EmitTextureCall(std::ostream& pre, std::ostream& out, ast::CallExpression* expr, - const semantic::TextureIntrinsicCall* sem); + const semantic::Intrinsic* intrinsic); /// Handles generating a call to data packing intrinsic /// @param pre the preamble of the expression stream /// @param out the output of the expression stream /// @param expr the call expression + /// @param intrinsic the semantic information for the texture intrinsic /// @returns true if the call expression is emitted bool EmitDataPackingCall(std::ostream& pre, std::ostream& out, - ast::CallExpression* expr); + ast::CallExpression* expr, + const semantic::Intrinsic* intrinsic); /// Handles a case statement /// @param out the output stream /// @param stmt the statement @@ -363,9 +365,9 @@ class GeneratorImpl { std::string generate_storage_buffer_index_expression(std::ostream& pre, ast::Expression* expr); /// Handles generating a builtin method name - /// @param call the semantic info for the intrinsic call + /// @param intrinsic the semantic info for the intrinsic /// @returns the name or "" if not valid - std::string generate_builtin_name(const semantic::IntrinsicCall* call); + std::string generate_builtin_name(const semantic::Intrinsic* intrinsic); /// Converts a builtin to an attribute name /// @param builtin the builtin to convert /// @returns the string name of the builtin or blank on error diff --git a/src/writer/hlsl/generator_impl_intrinsic_test.cc b/src/writer/hlsl/generator_impl_intrinsic_test.cc index 2370d844ef..2b32853765 100644 --- a/src/writer/hlsl/generator_impl_intrinsic_test.cc +++ b/src/writer/hlsl/generator_impl_intrinsic_test.cc @@ -175,7 +175,9 @@ TEST_P(HlslIntrinsicTest, Emit) { auto* sem = program->Sem().Get(call); ASSERT_NE(sem, nullptr); - auto* intrinsic = sem->As(); + auto* target = sem->Target(); + ASSERT_NE(target, nullptr); + auto* intrinsic = target->As(); ASSERT_NE(intrinsic, nullptr); EXPECT_EQ(gen.generate_builtin_name(intrinsic), param.hlsl_name); diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc index 002050122e..5c9c1def74 100644 --- a/src/writer/msl/generator_impl.cc +++ b/src/writer/msl/generator_impl.cc @@ -444,12 +444,12 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) { return 0; } - auto* call_sem = program_->Sem().Get(expr); - if (auto* sem = call_sem->As()) { - return EmitTextureCall(expr, sem); - } - if (auto* sem = call_sem->As()) { - if (sem->intrinsic() == semantic::IntrinsicType::kPack2x16Float) { + auto* call = program_->Sem().Get(expr); + if (auto* intrinsic = call->Target()->As()) { + if (intrinsic->IsTexture()) { + return EmitTextureCall(expr, intrinsic); + } + if (intrinsic->Type() == semantic::IntrinsicType::kPack2x16Float) { make_indent(); out_ << "as_type(half2("; if (!EmitExpression(expr->params()[0])) { @@ -458,7 +458,7 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) { out_ << "))"; return true; } - auto name = generate_builtin_name(sem); + auto name = generate_builtin_name(intrinsic); if (name.empty()) { return false; } @@ -568,18 +568,24 @@ bool GeneratorImpl::EmitCall(ast::CallExpression* expr) { } bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr, - const semantic::TextureIntrinsicCall* sem) { - auto* ident = expr->func()->As(); + const semantic::Intrinsic* intrinsic) { + using Usage = semantic::Parameter::Usage; - auto params = expr->params(); - auto& pidx = sem->Params().idx; - auto const kNotUsed = semantic::TextureIntrinsicCall::Parameters::kNotUsed; + auto parameters = intrinsic->Parameters(); + auto arguments = expr->params(); - assert(pidx.texture != kNotUsed); - auto* texture_type = - TypeOf(params[pidx.texture])->UnwrapAll()->As(); + // Returns the argument with the given usage + auto arg = [&](Usage usage) { + int idx = semantic::IndexOf(parameters, usage); + return (idx >= 0) ? arguments[idx] : nullptr; + }; - switch (sem->intrinsic()) { + auto* texture = arg(Usage::kTexture); + assert(texture); + + auto* texture_type = TypeOf(texture)->UnwrapAll()->As(); + + switch (intrinsic->Type()) { case semantic::IntrinsicType::kTextureDimensions: { std::vector dims; switch (texture_type->dim()) { @@ -606,12 +612,14 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr, } auto get_dim = [&](const char* name) { - if (!EmitExpression(params[pidx.texture])) { + if (!EmitExpression(texture)) { return false; } out_ << ".get_" << name << "("; - if (pidx.level != kNotUsed) { - out_ << pidx.level; + if (auto* level = arg(Usage::kLevel)) { + if (!EmitExpression(level)) { + return false; + } } out_ << ")"; return true; @@ -636,7 +644,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr, } case semantic::IntrinsicType::kTextureNumLayers: { out_ << "int("; - if (!EmitExpression(params[pidx.texture])) { + if (!EmitExpression(texture)) { return false; } out_ << ".get_array_size())"; @@ -644,7 +652,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr, } case semantic::IntrinsicType::kTextureNumLevels: { out_ << "int("; - if (!EmitExpression(params[pidx.texture])) { + if (!EmitExpression(texture)) { return false; } out_ << ".get_num_mip_levels())"; @@ -652,7 +660,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr, } case semantic::IntrinsicType::kTextureNumSamples: { out_ << "int("; - if (!EmitExpression(params[pidx.texture])) { + if (!EmitExpression(texture)) { return false; } out_ << ".get_num_samples())"; @@ -662,12 +670,12 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr, break; } - if (!EmitExpression(params[pidx.texture])) + if (!EmitExpression(texture)) return false; bool lod_param_is_named = true; - switch (sem->intrinsic()) { + switch (intrinsic->Type()) { case semantic::IntrinsicType::kTextureSample: case semantic::IntrinsicType::kTextureSampleBias: case semantic::IntrinsicType::kTextureSampleLevel: @@ -686,7 +694,7 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr, break; default: error_ = "Internal compiler error: Unhandled texture intrinsic '" + - program_->Symbols().NameFor(ident->symbol()) + "'"; + std::string(intrinsic->str()) + "'"; return false; } @@ -698,40 +706,38 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr, first_arg = false; }; - for (auto idx : {pidx.value, pidx.sampler, pidx.coords, pidx.array_index, - pidx.depth_ref, pidx.sample_index}) { - if (idx != kNotUsed) { + for (auto usage : + {Usage::kValue, Usage::kSampler, Usage::kCoords, Usage::kArrayIndex, + Usage::kDepthRef, Usage::kSampleIndex}) { + if (auto* e = arg(usage)) { maybe_write_comma(); - if (!EmitExpression(params[idx])) + if (!EmitExpression(e)) return false; } } - if (pidx.bias != kNotUsed) { + if (auto* bias = arg(Usage::kBias)) { maybe_write_comma(); out_ << "bias("; - if (!EmitExpression(params[pidx.bias])) { + if (!EmitExpression(bias)) { return false; } out_ << ")"; } - if (pidx.level != kNotUsed) { + if (auto* level = arg(Usage::kLevel)) { maybe_write_comma(); if (lod_param_is_named) { out_ << "level("; } - if (!EmitExpression(params[pidx.level])) { + if (!EmitExpression(level)) { return false; } if (lod_param_is_named) { out_ << ")"; } } - if (pidx.ddx != kNotUsed) { - auto dim = TypeOf(params[pidx.texture]) - ->UnwrapPtrIfNeeded() - ->As() - ->dim(); + if (auto* ddx = arg(Usage::kDdx)) { + auto dim = texture_type->dim(); switch (dim) { case type::TextureDimension::k2d: case type::TextureDimension::k2dArray: @@ -754,19 +760,19 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr, return false; } } - if (!EmitExpression(params[pidx.ddx])) { + if (!EmitExpression(ddx)) { return false; } out_ << ", "; - if (!EmitExpression(params[pidx.ddy])) { + if (!EmitExpression(arg(Usage::kDdy))) { return false; } out_ << ")"; } - if (pidx.offset != kNotUsed) { + if (auto* offset = arg(Usage::kOffset)) { maybe_write_comma(); - if (!EmitExpression(params[pidx.offset])) { + if (!EmitExpression(offset)) { return false; } } @@ -777,9 +783,9 @@ bool GeneratorImpl::EmitTextureCall(ast::CallExpression* expr, } std::string GeneratorImpl::generate_builtin_name( - const semantic::IntrinsicCall* call) { + const semantic::Intrinsic* intrinsic) { std::string out = "metal::"; - switch (call->intrinsic()) { + switch (intrinsic->Type()) { case semantic::IntrinsicType::kAcos: case semantic::IntrinsicType::kAll: case semantic::IntrinsicType::kAny: @@ -817,10 +823,10 @@ std::string GeneratorImpl::generate_builtin_name( case semantic::IntrinsicType::kTrunc: case semantic::IntrinsicType::kSign: case semantic::IntrinsicType::kClamp: - out += semantic::intrinsic::str(call->intrinsic()); + out += intrinsic->str(); break; case semantic::IntrinsicType::kAbs: - if (call->Type()->is_float_scalar_or_vector()) { + if (intrinsic->ReturnType()->is_float_scalar_or_vector()) { out += "fabs"; } else { out += "abs"; @@ -857,14 +863,14 @@ std::string GeneratorImpl::generate_builtin_name( out += "isnormal"; break; case semantic::IntrinsicType::kMax: - if (call->Type()->is_float_scalar_or_vector()) { + if (intrinsic->ReturnType()->is_float_scalar_or_vector()) { out += "fmax"; } else { out += "max"; } break; case semantic::IntrinsicType::kMin: - if (call->Type()->is_float_scalar_or_vector()) { + if (intrinsic->ReturnType()->is_float_scalar_or_vector()) { out += "fmin"; } else { out += "min"; @@ -895,8 +901,7 @@ std::string GeneratorImpl::generate_builtin_name( out += "rsqrt"; break; default: - error_ = "Unknown import method: " + - std::string(semantic::intrinsic::str(call->intrinsic())); + error_ = "Unknown import method: " + std::string(intrinsic->str()); return ""; } return out; diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h index 617061f303..ccbb0d17df 100644 --- a/src/writer/msl/generator_impl.h +++ b/src/writer/msl/generator_impl.h @@ -50,8 +50,8 @@ namespace tint { // Forward declarations namespace semantic { -class TextureIntrinsicCall; -class IntrinsicCall; +class Call; +class Intrinsic; } // namespace semantic namespace writer { @@ -121,10 +121,10 @@ class GeneratorImpl : public TextGenerator { /// Handles generating a call to a texture function (`textureSample`, /// `textureSampleGrad`, etc) /// @param expr the call expression - /// @param sem the semantic information for the texture intrinsic call + /// @param intrinsic the semantic information for the texture intrinsic /// @returns true if the call expression is emitted bool EmitTextureCall(ast::CallExpression* expr, - const semantic::TextureIntrinsicCall* sem); + const semantic::Intrinsic* intrinsic); /// Handles a case statement /// @param stmt the statement /// @returns true if the statement was emitted successfully @@ -258,9 +258,9 @@ class GeneratorImpl : public TextGenerator { /// @returns the name std::string generate_name(const std::string& prefix); /// Handles generating a builtin name - /// @param call the semantic info for the intrinsic call + /// @param intrinsic the semantic info for the intrinsic /// @returns the name or "" if not valid - std::string generate_builtin_name(const semantic::IntrinsicCall* call); + std::string generate_builtin_name(const semantic::Intrinsic* intrinsic); /// Checks if the global variable is in an input or output struct /// @param var the variable to check diff --git a/src/writer/msl/generator_impl_import_test.cc b/src/writer/msl/generator_impl_import_test.cc index 76440d0552..7b365f4d96 100644 --- a/src/writer/msl/generator_impl_import_test.cc +++ b/src/writer/msl/generator_impl_import_test.cc @@ -60,7 +60,9 @@ TEST_P(MslImportData_SingleParamTest, FloatScalar) { auto* sem = program->Sem().Get(call); ASSERT_NE(sem, nullptr); - auto* intrinsic = sem->As(); + auto* target = sem->Target(); + ASSERT_NE(target, nullptr); + auto* intrinsic = target->As(); ASSERT_NE(intrinsic, nullptr); ASSERT_EQ(gen.generate_builtin_name(intrinsic), diff --git a/src/writer/msl/generator_impl_intrinsic_test.cc b/src/writer/msl/generator_impl_intrinsic_test.cc index 7829636181..56b194afc3 100644 --- a/src/writer/msl/generator_impl_intrinsic_test.cc +++ b/src/writer/msl/generator_impl_intrinsic_test.cc @@ -182,7 +182,9 @@ TEST_P(MslIntrinsicTest, Emit) { auto* sem = program->Sem().Get(call); ASSERT_NE(sem, nullptr); - auto* intrinsic = sem->As(); + auto* target = sem->Target(); + ASSERT_NE(target, nullptr); + auto* intrinsic = target->As(); ASSERT_NE(intrinsic, nullptr); EXPECT_EQ(gen.generate_builtin_name(intrinsic), param.msl_name); diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc index ce32ebb14e..386f5a98f8 100644 --- a/src/writer/spirv/builder.cc +++ b/src/writer/spirv/builder.cc @@ -168,10 +168,10 @@ type::Matrix* GetNestedMatrixType(type::Type* type) { return type->As(); } -uint32_t intrinsic_to_glsl_method(type::Type* type, IntrinsicType intrinsic) { - switch (intrinsic) { +uint32_t intrinsic_to_glsl_method(const semantic::Intrinsic* intrinsic) { + switch (intrinsic->Type()) { case IntrinsicType::kAbs: - if (type->is_float_scalar_or_vector()) { + if (intrinsic->ReturnType()->is_float_scalar_or_vector()) { return GLSLstd450FAbs; } else { return GLSLstd450SAbs; @@ -187,9 +187,9 @@ uint32_t intrinsic_to_glsl_method(type::Type* type, IntrinsicType intrinsic) { case IntrinsicType::kCeil: return GLSLstd450Ceil; case IntrinsicType::kClamp: - if (type->is_float_scalar_or_vector()) { + if (intrinsic->ReturnType()->is_float_scalar_or_vector()) { return GLSLstd450NClamp; - } else if (type->is_unsigned_scalar_or_vector()) { + } else if (intrinsic->ReturnType()->is_unsigned_scalar_or_vector()) { return GLSLstd450UClamp; } else { return GLSLstd450SClamp; @@ -229,17 +229,17 @@ uint32_t intrinsic_to_glsl_method(type::Type* type, IntrinsicType intrinsic) { case IntrinsicType::kLog2: return GLSLstd450Log2; case IntrinsicType::kMax: - if (type->is_float_scalar_or_vector()) { + if (intrinsic->ReturnType()->is_float_scalar_or_vector()) { return GLSLstd450NMax; - } else if (type->is_unsigned_scalar_or_vector()) { + } else if (intrinsic->ReturnType()->is_unsigned_scalar_or_vector()) { return GLSLstd450UMax; } else { return GLSLstd450SMax; } case IntrinsicType::kMin: - if (type->is_float_scalar_or_vector()) { + if (intrinsic->ReturnType()->is_float_scalar_or_vector()) { return GLSLstd450NMin; - } else if (type->is_unsigned_scalar_or_vector()) { + } else if (intrinsic->ReturnType()->is_unsigned_scalar_or_vector()) { return GLSLstd450UMin; } else { return GLSLstd450SMin; @@ -1826,12 +1826,13 @@ uint32_t Builder::GenerateCallExpression(ast::CallExpression* expr) { return 0; } - auto* sem = builder_.Sem().Get(expr); - if (auto* intrinsic = sem->As()) { - return GenerateIntrinsic(ident, expr, intrinsic); + auto* call = builder_.Sem().Get(expr); + auto* target = call->Target(); + if (auto* intrinsic = target->As()) { + return GenerateIntrinsic(expr, intrinsic); } - auto type_id = GenerateTypeIfNeeded(sem->Type()); + auto type_id = GenerateTypeIfNeeded(target->ReturnType()); if (type_id == 0) { return 0; } @@ -1865,31 +1866,27 @@ uint32_t Builder::GenerateCallExpression(ast::CallExpression* expr) { return result_id; } -uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident, - ast::CallExpression* call, - const semantic::IntrinsicCall* sem) { +uint32_t Builder::GenerateIntrinsic(ast::CallExpression* call, + const semantic::Intrinsic* intrinsic) { auto result = result_op(); auto result_id = result.to_i(); - auto result_type_id = GenerateTypeIfNeeded(TypeOf(call)); + auto result_type_id = GenerateTypeIfNeeded(intrinsic->ReturnType()); if (result_type_id == 0) { return 0; } - auto intrinsic = sem->intrinsic(); - - if (semantic::intrinsic::IsFineDerivative(intrinsic) || - semantic::intrinsic::IsCoarseDerivative(intrinsic)) { + if (intrinsic->IsFineDerivative() || intrinsic->IsCoarseDerivative()) { push_capability(SpvCapabilityDerivativeControl); } - if (semantic::intrinsic::IsImageQueryIntrinsic(intrinsic)) { + if (intrinsic->IsImageQuery()) { push_capability(SpvCapabilityImageQuery); } - if (auto* tex_sem = sem->As()) { - if (!GenerateTextureIntrinsic(ident, call, tex_sem, - Operand::Int(result_type_id), result)) { + if (intrinsic->IsTexture()) { + if (!GenerateTextureIntrinsic(call, intrinsic, Operand::Int(result_type_id), + result)) { return 0; } return result_id; @@ -1898,97 +1895,118 @@ uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident, OperandList params = {Operand::Int(result_type_id), result}; spv::Op op = spv::Op::OpNop; - if (intrinsic == IntrinsicType::kAny) { - op = spv::Op::OpAny; - } else if (intrinsic == IntrinsicType::kAll) { - op = spv::Op::OpAll; - } else if (intrinsic == IntrinsicType::kArrayLength) { - if (call->params().empty()) { - error_ = "missing param for runtime array length"; - return 0; - } - auto* arg = call->params()[0]; + switch (intrinsic->Type()) { + case IntrinsicType::kAny: + op = spv::Op::OpAny; + break; + case IntrinsicType::kAll: + op = spv::Op::OpAll; + break; + case IntrinsicType::kArrayLength: { + if (call->params().empty()) { + error_ = "missing param for runtime array length"; + return 0; + } + auto* arg = call->params()[0]; - auto* accessor = arg->As(); - if (accessor == nullptr) { - error_ = "invalid expression for array length"; - return 0; - } + auto* accessor = arg->As(); + if (accessor == nullptr) { + error_ = "invalid expression for array length"; + return 0; + } - auto struct_id = GenerateExpression(accessor->structure()); - if (struct_id == 0) { - return 0; - } - params.push_back(Operand::Int(struct_id)); + auto struct_id = GenerateExpression(accessor->structure()); + if (struct_id == 0) { + return 0; + } + params.push_back(Operand::Int(struct_id)); - auto* type = TypeOf(accessor->structure())->UnwrapAll(); - if (!type->Is()) { - error_ = - "invalid type (" + type->type_name() + ") for runtime array length"; - return 0; - } - // Runtime array must be the last member in the structure - params.push_back(Operand::Int( - uint32_t(type->As()->impl()->members().size() - 1))); + auto* type = TypeOf(accessor->structure())->UnwrapAll(); + if (!type->Is()) { + error_ = + "invalid type (" + type->type_name() + ") for runtime array length"; + return 0; + } + // Runtime array must be the last member in the structure + params.push_back(Operand::Int( + uint32_t(type->As()->impl()->members().size() - 1))); - if (!push_function_inst(spv::Op::OpArrayLength, params)) { - return 0; + if (!push_function_inst(spv::Op::OpArrayLength, params)) { + return 0; + } + return result_id; } - return result_id; - } else if (intrinsic == IntrinsicType::kCountOneBits) { - op = spv::Op::OpBitCount; - } else if (intrinsic == IntrinsicType::kDot) { - op = spv::Op::OpDot; - } else if (intrinsic == IntrinsicType::kDpdx) { - op = spv::Op::OpDPdx; - } else if (intrinsic == IntrinsicType::kDpdxCoarse) { - op = spv::Op::OpDPdxCoarse; - } else if (intrinsic == IntrinsicType::kDpdxFine) { - op = spv::Op::OpDPdxFine; - } else if (intrinsic == IntrinsicType::kDpdy) { - op = spv::Op::OpDPdy; - } else if (intrinsic == IntrinsicType::kDpdyCoarse) { - op = spv::Op::OpDPdyCoarse; - } else if (intrinsic == IntrinsicType::kDpdyFine) { - op = spv::Op::OpDPdyFine; - } else if (intrinsic == IntrinsicType::kFwidth) { - op = spv::Op::OpFwidth; - } else if (intrinsic == IntrinsicType::kFwidthCoarse) { - op = spv::Op::OpFwidthCoarse; - } else if (intrinsic == IntrinsicType::kFwidthFine) { - op = spv::Op::OpFwidthFine; - } else if (intrinsic == IntrinsicType::kIsInf) { - op = spv::Op::OpIsInf; - } else if (intrinsic == IntrinsicType::kIsNan) { - op = spv::Op::OpIsNan; - } else if (intrinsic == IntrinsicType::kReverseBits) { - op = spv::Op::OpBitReverse; - } else if (intrinsic == IntrinsicType::kSelect) { - op = spv::Op::OpSelect; - } else { - GenerateGLSLstd450Import(); + case IntrinsicType::kCountOneBits: + op = spv::Op::OpBitCount; + break; + case IntrinsicType::kDot: + op = spv::Op::OpDot; + break; + case IntrinsicType::kDpdx: + op = spv::Op::OpDPdx; + break; + case IntrinsicType::kDpdxCoarse: + op = spv::Op::OpDPdxCoarse; + break; + case IntrinsicType::kDpdxFine: + op = spv::Op::OpDPdxFine; + break; + case IntrinsicType::kDpdy: + op = spv::Op::OpDPdy; + break; + case IntrinsicType::kDpdyCoarse: + op = spv::Op::OpDPdyCoarse; + break; + case IntrinsicType::kDpdyFine: + op = spv::Op::OpDPdyFine; + break; + case IntrinsicType::kFwidth: + op = spv::Op::OpFwidth; + break; + case IntrinsicType::kFwidthCoarse: + op = spv::Op::OpFwidthCoarse; + break; + case IntrinsicType::kFwidthFine: + op = spv::Op::OpFwidthFine; + break; + case IntrinsicType::kIsInf: + op = spv::Op::OpIsInf; + break; + case IntrinsicType::kIsNan: + op = spv::Op::OpIsNan; + break; + case IntrinsicType::kReverseBits: + op = spv::Op::OpBitReverse; + break; + case IntrinsicType::kSelect: + op = spv::Op::OpSelect; + break; + default: { + GenerateGLSLstd450Import(); - auto set_iter = import_name_to_id_.find(kGLSLstd450); - if (set_iter == import_name_to_id_.end()) { - error_ = std::string("unknown import ") + kGLSLstd450; - return 0; + auto set_iter = import_name_to_id_.find(kGLSLstd450); + if (set_iter == import_name_to_id_.end()) { + error_ = std::string("unknown import ") + kGLSLstd450; + return 0; + } + auto set_id = set_iter->second; + auto inst_id = intrinsic_to_glsl_method(intrinsic); + if (inst_id == 0) { + error_ = "unknown method " + std::string(intrinsic->str()); + return 0; + } + + params.push_back(Operand::Int(set_id)); + params.push_back(Operand::Int(inst_id)); + + op = spv::Op::OpExtInst; + break; } - auto set_id = set_iter->second; - auto inst_id = intrinsic_to_glsl_method(sem->Type(), sem->intrinsic()); - if (inst_id == 0) { - error_ = "unknown method " + builder_.Symbols().NameFor(ident->symbol()); - return 0; - } - - params.push_back(Operand::Int(set_id)); - params.push_back(Operand::Int(inst_id)); - - op = spv::Op::OpExtInst; } if (op == spv::Op::OpNop) { - error_ = "unable to determine operator for: " + - builder_.Symbols().NameFor(ident->symbol()); + error_ = + "unable to determine operator for: " + std::string(intrinsic->str()); return 0; } @@ -2009,32 +2027,46 @@ uint32_t Builder::GenerateIntrinsic(ast::IdentifierExpression* ident, return result_id; } -bool Builder::GenerateTextureIntrinsic( - ast::IdentifierExpression* ident, - ast::CallExpression* call, - const semantic::TextureIntrinsicCall* sem, - Operand result_type, - Operand result_id) { - auto& pidx = sem->Params().idx; - auto const kNotUsed = semantic::TextureIntrinsicCall::Parameters::kNotUsed; +bool Builder::GenerateTextureIntrinsic(ast::CallExpression* call, + const semantic::Intrinsic* intrinsic, + Operand result_type, + Operand result_id) { + using Usage = semantic::Parameter::Usage; - assert(pidx.texture != kNotUsed); - auto* texture_type = - TypeOf(call->params()[pidx.texture])->UnwrapAll()->As(); + auto parameters = intrinsic->Parameters(); + auto arguments = call->params(); - auto op = spv::Op::OpNop; - - auto gen_param = [&](size_t idx) { - auto* p = call->params()[idx]; - auto val_id = GenerateExpression(p); + // Generates the given expression, returning the operand ID + auto gen = [&](ast::Expression* expr) { + auto val_id = GenerateExpression(expr); if (val_id == 0) { return Operand::Int(0); } - val_id = GenerateLoadIfNeeded(TypeOf(p), val_id); + val_id = GenerateLoadIfNeeded(TypeOf(expr), val_id); return Operand::Int(val_id); }; + // Returns the argument with the given usage + auto arg = [&](Usage usage) { + int idx = semantic::IndexOf(parameters, usage); + return (idx >= 0) ? arguments[idx] : nullptr; + }; + + // Generates the argument with the given usage, returning the operand ID + auto gen_arg = [&](Usage usage) { + auto* argument = arg(usage); + assert(argument); + return gen(argument); + }; + + auto* texture = arg(Usage::kTexture); + assert(texture); + + auto* texture_type = TypeOf(texture)->UnwrapAll()->As(); + + auto op = spv::Op::OpNop; + // Custom function to call after the texture-intrinsic op has been generated. std::function post_emission = [] { return true; }; @@ -2133,28 +2165,23 @@ bool Builder::GenerateTextureIntrinsic( }; auto append_coords_to_spirv_params = [&]() -> bool { - if (pidx.array_index != kNotUsed) { + if (auto* array_index = arg(Usage::kArrayIndex)) { // Array index needs to be appended to the coordinates. - auto* param_coords = call->params()[pidx.coords]; - auto* param_array_index = call->params()[pidx.array_index]; - - auto* packed = AppendVector(&builder_, param_coords, param_array_index); + auto* packed = AppendVector(&builder_, arg(Usage::kCoords), array_index); auto param = GenerateTypeConstructorExpression(packed, false); if (param == 0) { return false; } spirv_params.emplace_back(Operand::Int(param)); } else { - spirv_params.emplace_back(gen_param(pidx.coords)); // coordinates + spirv_params.emplace_back(gen_arg(Usage::kCoords)); // coordinates } return true; }; auto append_image_and_coords_to_spirv_params = [&]() -> bool { - assert(pidx.sampler != kNotUsed); - assert(pidx.texture != kNotUsed); - auto sampler_param = gen_param(pidx.sampler); - auto texture_param = gen_param(pidx.texture); + auto sampler_param = gen_arg(Usage::kSampler); + auto texture_param = gen_arg(Usage::kTexture); auto sampled_image = GenerateSampledImage(texture_type, texture_param, sampler_param); @@ -2163,7 +2190,7 @@ bool Builder::GenerateTextureIntrinsic( return append_coords_to_spirv_params(); }; - switch (sem->intrinsic()) { + switch (intrinsic->Type()) { case IntrinsicType::kTextureDimensions: { // Number of returned elements from OpImageQuerySize[Lod] may not match // those of textureDimensions(). @@ -2205,13 +2232,13 @@ bool Builder::GenerateTextureIntrinsic( return false; } - spirv_params.emplace_back(gen_param(pidx.texture)); + spirv_params.emplace_back(gen_arg(Usage::kTexture)); if (texture_type->Is() || texture_type->Is()) { op = spv::Op::OpImageQuerySize; - } else if (pidx.level != kNotUsed) { + } else if (auto* level = arg(Usage::kLevel)) { op = spv::Op::OpImageQuerySizeLod; - spirv_params.emplace_back(gen_param(pidx.level)); + spirv_params.emplace_back(gen(level)); } else { ast::SintLiteral i32_0(Source{}, builder_.create(), 0); op = spv::Op::OpImageQuerySizeLod; @@ -2242,7 +2269,7 @@ bool Builder::GenerateTextureIntrinsic( return false; } - spirv_params.emplace_back(gen_param(pidx.texture)); + spirv_params.emplace_back(gen_arg(Usage::kTexture)); if (texture_type->Is() || texture_type->Is()) { @@ -2258,43 +2285,43 @@ bool Builder::GenerateTextureIntrinsic( case IntrinsicType::kTextureNumLevels: { op = spv::Op::OpImageQueryLevels; append_result_type_and_id_to_spirv_params(); - spirv_params.emplace_back(gen_param(pidx.texture)); + spirv_params.emplace_back(gen_arg(Usage::kTexture)); break; } case IntrinsicType::kTextureNumSamples: { op = spv::Op::OpImageQuerySamples; append_result_type_and_id_to_spirv_params(); - spirv_params.emplace_back(gen_param(pidx.texture)); + spirv_params.emplace_back(gen_arg(Usage::kTexture)); break; } case IntrinsicType::kTextureLoad: { op = texture_type->Is() ? spv::Op::OpImageRead : spv::Op::OpImageFetch; append_result_type_and_id_to_spirv_params_for_read(); - spirv_params.emplace_back(gen_param(pidx.texture)); + spirv_params.emplace_back(gen_arg(Usage::kTexture)); if (!append_coords_to_spirv_params()) { return false; } - if (pidx.level != kNotUsed) { + if (auto* level = arg(Usage::kLevel)) { image_operands.emplace_back( - ImageOperand{SpvImageOperandsLodMask, gen_param(pidx.level)}); + ImageOperand{SpvImageOperandsLodMask, gen(level)}); } - if (pidx.sample_index != kNotUsed) { - image_operands.emplace_back(ImageOperand{SpvImageOperandsSampleMask, - gen_param(pidx.sample_index)}); + if (auto* sample_index = arg(Usage::kSampleIndex)) { + image_operands.emplace_back( + ImageOperand{SpvImageOperandsSampleMask, gen(sample_index)}); } break; } case IntrinsicType::kTextureStore: { op = spv::Op::OpImageWrite; - spirv_params.emplace_back(gen_param(pidx.texture)); + spirv_params.emplace_back(gen_arg(Usage::kTexture)); if (!append_coords_to_spirv_params()) { return false; } - spirv_params.emplace_back(gen_param(pidx.value)); + spirv_params.emplace_back(gen_arg(Usage::kValue)); break; } case IntrinsicType::kTextureSample: { @@ -2311,9 +2338,8 @@ bool Builder::GenerateTextureIntrinsic( if (!append_image_and_coords_to_spirv_params()) { return false; } - assert(pidx.bias != kNotUsed); image_operands.emplace_back( - ImageOperand{SpvImageOperandsBiasMask, gen_param(pidx.bias)}); + ImageOperand{SpvImageOperandsBiasMask, gen_arg(Usage::kBias)}); break; } case IntrinsicType::kTextureSampleLevel: { @@ -2322,20 +2348,19 @@ bool Builder::GenerateTextureIntrinsic( if (!append_image_and_coords_to_spirv_params()) { return false; } - assert(pidx.level != kNotUsed); auto level = Operand::Int(0); - if (TypeOf(call->params()[pidx.level])->Is()) { + if (TypeOf(arg(Usage::kLevel))->Is()) { // Depth textures have i32 parameters for the level, but SPIR-V expects // F32. Cast. auto* f32 = builder_.create(); ast::TypeConstructorExpression cast(Source{}, f32, - {call->params()[pidx.level]}); + {arg(Usage::kLevel)}); level = Operand::Int(GenerateExpression(&cast)); if (level.to_i() == 0) { return false; } } else { - level = gen_param(pidx.level); + level = gen_arg(Usage::kLevel); } image_operands.emplace_back(ImageOperand{SpvImageOperandsLodMask, level}); break; @@ -2346,12 +2371,10 @@ bool Builder::GenerateTextureIntrinsic( if (!append_image_and_coords_to_spirv_params()) { return false; } - assert(pidx.ddx != kNotUsed); - assert(pidx.ddy != kNotUsed); image_operands.emplace_back( - ImageOperand{SpvImageOperandsGradMask, gen_param(pidx.ddx)}); + ImageOperand{SpvImageOperandsGradMask, gen_arg(Usage::kDdx)}); image_operands.emplace_back( - ImageOperand{SpvImageOperandsGradMask, gen_param(pidx.ddy)}); + ImageOperand{SpvImageOperandsGradMask, gen_arg(Usage::kDdy)}); break; } case IntrinsicType::kTextureSampleCompare: { @@ -2360,8 +2383,7 @@ bool Builder::GenerateTextureIntrinsic( if (!append_image_and_coords_to_spirv_params()) { return false; } - assert(pidx.depth_ref != kNotUsed); - spirv_params.emplace_back(gen_param(pidx.depth_ref)); + spirv_params.emplace_back(gen_arg(Usage::kDepthRef)); type::F32 f32; ast::FloatLiteral float_0(Source{}, &f32, 0.0); @@ -2374,9 +2396,9 @@ bool Builder::GenerateTextureIntrinsic( break; // unreachable } - if (pidx.offset != kNotUsed) { + if (auto* offset = arg(Usage::kOffset)) { image_operands.emplace_back( - ImageOperand{SpvImageOperandsConstOffsetMask, gen_param(pidx.offset)}); + ImageOperand{SpvImageOperandsConstOffsetMask, gen(offset)}); } if (!image_operands.empty()) { @@ -2393,8 +2415,8 @@ bool Builder::GenerateTextureIntrinsic( } if (op == spv::Op::OpNop) { - error_ = "unable to determine operator for: " + - builder_.Symbols().NameFor(ident->symbol()); + error_ = + "unable to determine operator for: " + std::string(intrinsic->str()); return false; } diff --git a/src/writer/spirv/builder.h b/src/writer/spirv/builder.h index 3ab7479142..b1d38b3c8d 100644 --- a/src/writer/spirv/builder.h +++ b/src/writer/spirv/builder.h @@ -60,8 +60,7 @@ namespace tint { // Forward declarations namespace semantic { -class TextureIntrinsicCall; -class IntrinsicCall; +class Call; } // namespace semantic namespace writer { @@ -359,25 +358,21 @@ class Builder { /// @returns the expression ID on success or 0 otherwise uint32_t GenerateCallExpression(ast::CallExpression* expr); /// Generates an intrinsic call - /// @param ident the intrinsic expression /// @param call the call expression - /// @param sem the semantic information for the intrinsic call + /// @param intrinsic the semantic information for the intrinsic /// @returns the expression ID on success or 0 otherwise - uint32_t GenerateIntrinsic(ast::IdentifierExpression* ident, - ast::CallExpression* call, - const semantic::IntrinsicCall* sem); + uint32_t GenerateIntrinsic(ast::CallExpression* call, + const semantic::Intrinsic* intrinsic); /// Generates a texture intrinsic call. Emits an error and returns false if /// we're currently outside a function. - /// @param ident the texture intrinsic /// @param call the call expression - /// @param sem the semantic information for the texture intrinsic call + /// @param intrinsic the semantic information for the texture intrinsic /// @param result_type result type operand of the texture instruction /// @param result_id result identifier operand of the texture instruction /// parameters /// @returns true on success - bool GenerateTextureIntrinsic(ast::IdentifierExpression* ident, - ast::CallExpression* call, - const semantic::TextureIntrinsicCall* sem, + bool GenerateTextureIntrinsic(ast::CallExpression* call, + const semantic::Intrinsic* intrinsic, spirv::Operand result_type, spirv::Operand result_id); /// Generates a sampled image