Add semantic::Intrinsic
semantic::Intrinsic derives from semantic::CallTarget, which can be obtained from the Target() accessor on the CallExpression.
Flesh out semantic::Parameter to contain a `Usage` - extra metadata for the parameter.
The information in `Intrinsic` is enough to remove the `semantic::IntrinsicCall` and `semantic::TextureIntrinsicCall` types.
Change-Id: Ida9c193674ad8605d8f12f6a1d27f38c7d008434
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/40503
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/semantic/call.h b/src/semantic/call.h
index 6351421..0ef874f 100644
--- a/src/semantic/call.h
+++ b/src/semantic/call.h
@@ -26,96 +26,17 @@
class Call : public Castable<Call, Expression> {
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<IntrinsicCall, Call> {
- 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<TextureIntrinsicCall, IntrinsicCall> {
- public:
- /// Parameters describes the parameters for the texture function.
- struct Parameters {
- /// kNotUsed is the constant that indicates the given parameter is not part
- /// of the texture function signature.
- static constexpr const size_t kNotUsed = ~static_cast<size_t>(0u);
- /// Index holds each of the possible parameter indices. If a parameter index
- /// is equal to `kNotUsed` then this parameter is not used by the function.
- struct Index {
- /// Constructor
- Index();
- /// Copy constructor
- Index(const Index&);
- /// `array_index` parameter index.
- size_t array_index = kNotUsed;
- /// `bias` parameter index.
- size_t bias = kNotUsed;
- /// `coords` parameter index.
- size_t coords = kNotUsed;
- /// `depth_ref` parameter index.
- size_t depth_ref = kNotUsed;
- /// `ddx` parameter index.
- size_t ddx = kNotUsed;
- /// `ddy` parameter index.
- size_t ddy = kNotUsed;
- /// `level` parameter index.
- size_t level = kNotUsed;
- /// `offset` parameter index.
- size_t offset = kNotUsed;
- /// `sampler` parameter index.
- size_t sampler = kNotUsed;
- /// `sample_index` parameter index.
- size_t sample_index = kNotUsed;
- /// `texture` parameter index.
- size_t texture = kNotUsed;
- /// `value` parameter index.
- size_t value = kNotUsed;
- };
- /// The indices of all possible parameters.
- Index idx;
- /// Total number of parameters.
- size_t count = 0;
- };
-
- /// Constructor
- /// @param return_type the return type of the call
- /// @param intrinsic the call target intrinsic
- /// @param params the overload parameter info
- TextureIntrinsicCall(type::Type* return_type,
- 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 ddb6aac..43cbc0e 100644
--- a/src/semantic/call_target.h
+++ b/src/semantic/call_target.h
@@ -32,18 +32,52 @@
/// 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<Parameter>;
+/// @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<CallTarget, Node> {
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 @@
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 f006f31..e2c1902 100644
--- a/src/semantic/intrinsic.h
+++ b/src/semantic/intrinsic.h
@@ -17,6 +17,8 @@
#include <ostream>
+#include "src/semantic/call_target.h"
+
namespace tint {
namespace semantic {
@@ -102,39 +104,37 @@
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 @@
/// @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<Intrinsic, CallTarget> {
+ 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);
-} // namespace intrinsic
+ /// 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 semantic
} // namespace tint
diff --git a/src/semantic/sem_call.cc b/src/semantic/sem_call.cc
index 777ad51..632f3a0 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 3b79a55..ad7f300 100644
--- a/src/semantic/sem_call_target.cc
+++ b/src/semantic/sem_call_target.cc
@@ -21,10 +21,50 @@
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<int>(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 "<unknown>";
+ }
+}
} // namespace semantic
} // namespace tint
diff --git a/src/semantic/sem_function.cc b/src/semantic/sem_function.cc
index 8c8d16c..db53339 100644
--- a/src/semantic/sem_function.cc
+++ b/src/semantic/sem_function.cc
@@ -37,7 +37,7 @@
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 @@
std::vector<const Variable*> referenced_module_vars,
std::vector<const Variable*> local_referenced_module_vars,
std::vector<Symbol> 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 b2eea27..e12c1b4 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 @@
return "<unknown>";
}
-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 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 @@
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 6b47ffa..7ea3c09 100644
--- a/src/type_determiner.cc
+++ b/src/type_determiner.cc
@@ -419,9 +419,9 @@
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 @@
}
auto* function = iter->second;
- auto* return_ty = function->declaration->return_type();
- auto* sem = builder_->create<semantic::Call>(return_ty);
- builder_->Sem().Add(call, sem);
+ function_calls_.emplace(call, function);
}
return true;
@@ -548,14 +546,26 @@
} // namespace
bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
- IntrinsicType intrinsic) {
- auto create_sem = [&](type::Type* result) {
- auto* sem = builder_->create<semantic::IntrinsicCall>(result, intrinsic);
- builder_->Sem().Add(call, sem);
+ IntrinsicType intrinsic_type) {
+ using Parameter = semantic::Parameter;
+ using Parameters = semantic::Parameters;
+ using Usage = Parameter::Usage;
+
+ std::vector<type::Type*> 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<semantic::Intrinsic>(
+ intrinsic_type, return_type, params);
+ builder_->Sem().Add(call, builder_->create<semantic::Call>(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 @@
}
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<type::Texture>()) {
+ auto& ty = builder_->ty;
+
+ auto* texture = arg_tys[0]->UnwrapAll()->As<type::Texture>();
+ if (!texture) {
set_error(call->source(), "invalid first argument for " + name);
return false;
}
- type::Texture* texture =
- TypeOf(texture_param)->UnwrapAll()->As<type::Texture>();
bool is_array = type::IsTextureArray(texture->dim());
bool is_multisampled = texture->Is<type::MultisampledTexture>();
- 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<int>(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<type::I32>();
switch (texture->dim()) {
@@ -748,16 +762,15 @@
}
}
- auto* sem = builder_->create<semantic::TextureIntrinsicCall>(
- return_type, intrinsic, param);
- builder_->Sem().Add(call, sem);
-
+ auto* intrinsic = builder_->create<semantic::Intrinsic>(
+ intrinsic_type, return_type, params);
+ builder_->Sem().Add(call, builder_->create<semantic::Call>(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 @@
}
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 @@
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 @@
return out;
};
+ std::unordered_map<FunctionInfo*, semantic::Function*> func_info_to_sem_func;
for (auto it : function_to_info_) {
auto* func = it.first;
auto* info = it.second;
- sem.Add(func,
- builder_->create<semantic::Function>(
- info->declaration, remap_vars(info->referenced_module_vars),
- remap_vars(info->local_referenced_module_vars),
- info->ancestor_entry_points));
+ auto* sem_func = builder_->create<semantic::Function>(
+ 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<semantic::Call>(sem_func));
}
}
diff --git a/src/type_determiner.h b/src/type_determiner.h
index c8dc34d..9928b90 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 @@
/// @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 <typename T>
@@ -177,7 +178,7 @@
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 @@
/// @param type the resolved type
void SetType(ast::Expression* expr, type::Type* type) const;
- ProgramBuilder* builder_;
+ ProgramBuilder* const builder_;
std::string error_;
ScopeStack<VariableInfo*> variable_stack_;
std::unordered_map<Symbol, FunctionInfo*> symbol_to_function_;
std::unordered_map<ast::Function*, FunctionInfo*> function_to_info_;
std::unordered_map<ast::Variable*, VariableInfo*> variable_to_info_;
+ std::unordered_map<ast::CallExpression*, FunctionInfo*> function_calls_;
FunctionInfo* current_function_ = nullptr;
BlockAllocator<VariableInfo> variable_infos_;
BlockAllocator<FunctionInfo> function_infos_;
diff --git a/src/type_determiner_test.cc b/src/type_determiner_test.cc
index 72f9e59..4b43554 100644
--- a/src/type_determiner_test.cc
+++ b/src/type_determiner_test.cc
@@ -1637,7 +1637,7 @@
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 @@
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 @@
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<Parameter> list;
- auto maybe_add_param = [&list](size_t idx, const char* name) {
- if (idx !=
- semantic::TextureIntrinsicCall::Parameters::Parameters::kNotUsed) {
- list.emplace_back(Parameter{idx, name});
- }
- };
- maybe_add_param(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 @@
}
}
- auto* sem = Sem().Get(call);
- ASSERT_NE(sem, nullptr);
- auto* intrinsic = sem->As<semantic::TextureIntrinsicCall>();
- 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 ecad8ef..019423e 100644
--- a/src/validator/validator_impl.cc
+++ b/src/validator/validator_impl.cc
@@ -644,24 +644,23 @@
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<semantic::IntrinsicCall>()) {
+ if (auto* intrinsic = call->Target()->As<semantic::Intrinsic>()) {
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 @@
}
}
- if (semantic::intrinsic::IsDataPackingIntrinsic(data->intrinsic)) {
+ if (semantic::IsDataPackingIntrinsic(data->intrinsic)) {
if (!program_->TypeOf(expr)->Is<type::U32>()) {
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 2a4d304..9ac33c8 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -542,22 +542,22 @@
return 0;
}
- auto* call_sem = builder_.Sem().Get(expr);
- if (auto* sem = call_sem->As<semantic::TextureIntrinsicCall>()) {
- return EmitTextureCall(pre, out, expr, sem);
- }
- if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) {
+ auto* call = builder_.Sem().Get(expr);
+ if (auto* intrinsic = call->Target()->As<semantic::Intrinsic>()) {
+ 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::EmitDataPackingCall(std::ostream& pre,
std::ostream& out,
- ast::CallExpression* expr) {
- auto* ident = builder_.Sem().Get(expr)->As<semantic::IntrinsicCall>();
-
+ 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 @@
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::EmitTextureCall(std::ostream& pre,
std::ostream& out,
ast::CallExpression* expr,
- const semantic::TextureIntrinsicCall* sem) {
- auto* ident = expr->func()->As<ast::IdentifierExpression>();
+ 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();
- auto* texture = params[pidx.texture];
+ // 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_type = TypeOf(texture)->UnwrapAll()->As<type::Texture>();
- switch (sem->intrinsic()) {
+ switch (intrinsic->Type()) {
case semantic::IntrinsicType::kTextureDimensions:
case semantic::IntrinsicType::kTextureNumLayers:
case semantic::IntrinsicType::kTextureNumLevels:
@@ -718,7 +724,7 @@
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 @@
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 pack_mip_in_coords = false;
- switch (sem->intrinsic()) {
+ switch (intrinsic->Type()) {
case semantic::IntrinsicType::kTextureSample:
out << ".Sample(";
break;
@@ -886,17 +895,18 @@
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<type::I32>();
@@ -906,10 +916,9 @@
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 @@
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 @@
} // 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 @@
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 @@
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 22b8be0..805a0ce 100644
--- a/src/writer/hlsl/generator_impl.h
+++ b/src/writer/hlsl/generator_impl.h
@@ -48,8 +48,8 @@
// Forward declarations
namespace semantic {
-class TextureIntrinsicCall;
-class IntrinsicCall;
+class Call;
+class Intrinsic;
} // namespace semantic
namespace writer {
@@ -153,20 +153,22 @@
/// @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 @@
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 2370d84..2b32853 100644
--- a/src/writer/hlsl/generator_impl_intrinsic_test.cc
+++ b/src/writer/hlsl/generator_impl_intrinsic_test.cc
@@ -175,7 +175,9 @@
auto* sem = program->Sem().Get(call);
ASSERT_NE(sem, nullptr);
- auto* intrinsic = sem->As<semantic::IntrinsicCall>();
+ auto* target = sem->Target();
+ ASSERT_NE(target, nullptr);
+ auto* intrinsic = target->As<semantic::Intrinsic>();
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 0020501..5c9c1de 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -444,12 +444,12 @@
return 0;
}
- auto* call_sem = program_->Sem().Get(expr);
- if (auto* sem = call_sem->As<semantic::TextureIntrinsicCall>()) {
- return EmitTextureCall(expr, sem);
- }
- if (auto* sem = call_sem->As<semantic::IntrinsicCall>()) {
- if (sem->intrinsic() == semantic::IntrinsicType::kPack2x16Float) {
+ auto* call = program_->Sem().Get(expr);
+ if (auto* intrinsic = call->Target()->As<semantic::Intrinsic>()) {
+ if (intrinsic->IsTexture()) {
+ return EmitTextureCall(expr, intrinsic);
+ }
+ if (intrinsic->Type() == semantic::IntrinsicType::kPack2x16Float) {
make_indent();
out_ << "as_type<uint>(half2(";
if (!EmitExpression(expr->params()[0])) {
@@ -458,7 +458,7 @@
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::EmitTextureCall(ast::CallExpression* expr,
- const semantic::TextureIntrinsicCall* sem) {
- auto* ident = expr->func()->As<ast::IdentifierExpression>();
+ 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<type::Texture>();
+ // 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<type::Texture>();
+
+ switch (intrinsic->Type()) {
case semantic::IntrinsicType::kTextureDimensions: {
std::vector<const char*> dims;
switch (texture_type->dim()) {
@@ -606,12 +612,14 @@
}
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 @@
}
case semantic::IntrinsicType::kTextureNumLayers: {
out_ << "int(";
- if (!EmitExpression(params[pidx.texture])) {
+ if (!EmitExpression(texture)) {
return false;
}
out_ << ".get_array_size())";
@@ -644,7 +652,7 @@
}
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 @@
}
case semantic::IntrinsicType::kTextureNumSamples: {
out_ << "int(";
- if (!EmitExpression(params[pidx.texture])) {
+ if (!EmitExpression(texture)) {
return false;
}
out_ << ".get_num_samples())";
@@ -662,12 +670,12 @@
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 @@
break;
default:
error_ = "Internal compiler error: Unhandled texture intrinsic '" +
- program_->Symbols().NameFor(ident->symbol()) + "'";
+ std::string(intrinsic->str()) + "'";
return false;
}
@@ -698,40 +706,38 @@
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<type::Texture>()
- ->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 @@
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 @@
}
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 @@
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 @@
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 @@
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 617061f..ccbb0d1 100644
--- a/src/writer/msl/generator_impl.h
+++ b/src/writer/msl/generator_impl.h
@@ -50,8 +50,8 @@
// Forward declarations
namespace semantic {
-class TextureIntrinsicCall;
-class IntrinsicCall;
+class Call;
+class Intrinsic;
} // namespace semantic
namespace writer {
@@ -121,10 +121,10 @@
/// Handles generating a call to a texture function (`textureSample`,
/// `textureSampleGrad`, etc)
/// @param expr the call expression
- /// @param sem the semantic information for the texture intrinsic call
+ /// @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 @@
/// @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 76440d0..7b365f4 100644
--- a/src/writer/msl/generator_impl_import_test.cc
+++ b/src/writer/msl/generator_impl_import_test.cc
@@ -60,7 +60,9 @@
auto* sem = program->Sem().Get(call);
ASSERT_NE(sem, nullptr);
- auto* intrinsic = sem->As<semantic::IntrinsicCall>();
+ auto* target = sem->Target();
+ ASSERT_NE(target, nullptr);
+ auto* intrinsic = target->As<semantic::Intrinsic>();
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 7829636..56b194a 100644
--- a/src/writer/msl/generator_impl_intrinsic_test.cc
+++ b/src/writer/msl/generator_impl_intrinsic_test.cc
@@ -182,7 +182,9 @@
auto* sem = program->Sem().Get(call);
ASSERT_NE(sem, nullptr);
- auto* intrinsic = sem->As<semantic::IntrinsicCall>();
+ auto* target = sem->Target();
+ ASSERT_NE(target, nullptr);
+ auto* intrinsic = target->As<semantic::Intrinsic>();
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 ce32ebb..386f5a9 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -168,10 +168,10 @@
return type->As<type::Matrix>();
}
-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 @@
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 @@
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 @@
return 0;
}
- auto* sem = builder_.Sem().Get(expr);
- if (auto* intrinsic = sem->As<semantic::IntrinsicCall>()) {
- return GenerateIntrinsic(ident, expr, intrinsic);
+ auto* call = builder_.Sem().Get(expr);
+ auto* target = call->Target();
+ if (auto* intrinsic = target->As<semantic::Intrinsic>()) {
+ 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 @@
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<semantic::TextureIntrinsicCall>()) {
- 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 @@
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<ast::MemberAccessorExpression>();
- if (accessor == nullptr) {
- error_ = "invalid expression for array length";
- return 0;
- }
+ auto* accessor = arg->As<ast::MemberAccessorExpression>();
+ 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<type::Struct>()) {
- 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<type::Struct>()->impl()->members().size() - 1)));
+ auto* type = TypeOf(accessor->structure())->UnwrapAll();
+ if (!type->Is<type::Struct>()) {
+ 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<type::Struct>()->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 @@
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<type::Texture>();
+ 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<type::Texture>();
+
+ auto op = spv::Op::OpNop;
+
// Custom function to call after the texture-intrinsic op has been generated.
std::function<bool()> post_emission = [] { return true; };
@@ -2133,28 +2165,23 @@
};
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 @@
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 @@
return false;
}
- spirv_params.emplace_back(gen_param(pidx.texture));
+ spirv_params.emplace_back(gen_arg(Usage::kTexture));
if (texture_type->Is<type::MultisampledTexture>() ||
texture_type->Is<type::StorageTexture>()) {
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<type::I32>(), 0);
op = spv::Op::OpImageQuerySizeLod;
@@ -2242,7 +2269,7 @@
return false;
}
- spirv_params.emplace_back(gen_param(pidx.texture));
+ spirv_params.emplace_back(gen_arg(Usage::kTexture));
if (texture_type->Is<type::MultisampledTexture>() ||
texture_type->Is<type::StorageTexture>()) {
@@ -2258,43 +2285,43 @@
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<type::StorageTexture>() ? 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 @@
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 @@
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<type::I32>()) {
+ if (TypeOf(arg(Usage::kLevel))->Is<type::I32>()) {
// Depth textures have i32 parameters for the level, but SPIR-V expects
// F32. Cast.
auto* f32 = builder_.create<type::F32>();
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 @@
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 @@
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 @@
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 @@
}
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 3ab7479..b1d38b3 100644
--- a/src/writer/spirv/builder.h
+++ b/src/writer/spirv/builder.h
@@ -60,8 +60,7 @@
// Forward declarations
namespace semantic {
-class TextureIntrinsicCall;
-class IntrinsicCall;
+class Call;
} // namespace semantic
namespace writer {
@@ -359,25 +358,21 @@
/// @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
