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

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>
This commit is contained in:
Ben Clayton 2021-02-08 22:31:44 +00:00 committed by Commit Bot service account
parent bab3197fda
commit 316f9f6b6d
20 changed files with 659 additions and 562 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

89
src/semantic/call.h
View File

@ -26,96 +26,17 @@ namespace semantic {
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

41
src/semantic/call_target.h
View File

@ -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<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 @@ class CallTarget : public Castable<CallTarget, Node> {
const Parameters& Parameters() const { return parameters_; }
private:
semantic::Parameters parameters_;
type::Type* const return_type_;
semantic::Parameters const parameters_;
};
} // namespace semantic

81
src/semantic/intrinsic.h
View File

@ -17,6 +17,8 @@
#include <ostream>
#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<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);
/// 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

20
src/semantic/sem_call.cc
View File

@ -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

44
src/semantic/sem_call_target.cc
View File

@ -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<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

4
src/semantic/sem_function.cc
View File

@ -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<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)) {}

56
src/semantic/sem_intrinsic.cc
View File

@ -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 "<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 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

198
src/type_determiner.cc
View File

@ -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<semantic::Call>(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<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 @@ 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<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 @@ bool TypeDeterminer::DetermineIntrinsicCall(ast::CallExpression* call,
}
}
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 @@ 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<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));
}
}

12
src/type_determiner.h
View File

@ -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 <typename T>
@ -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<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_;

47
src/type_determiner_test.cc
View File

@ -1637,7 +1637,7 @@ using IntrinsicDataTest = TypeDeterminerTestWithParam<IntrinsicData>;
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<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 @@ TEST_P(TypeDeterminerTextureIntrinsicTest, Call) {
}
}
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);
}

13
src/validator/validator_impl.cc
View File

@ -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<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 @@ bool ValidatorImpl::ValidateCallExpr(const ast::CallExpression* expr) {
}
}
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 +

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

@ -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<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::EmitCall(std::ostream& pre,
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 @@ 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<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();
// 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<type::Texture>();
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<type::I32>();
@ -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 "";
}

16
src/writer/hlsl/generator_impl.h
View File

@ -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

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

@ -175,7 +175,9 @@ TEST_P(HlslIntrinsicTest, Emit) {
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);

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

@ -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<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 @@ 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<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 @@ 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<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 @@ 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;

12
src/writer/msl/generator_impl.h
View File

@ -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

4
src/writer/msl/generator_impl_import_test.cc
View File

@ -60,7 +60,9 @@ TEST_P(MslImportData_SingleParamTest, FloatScalar) {
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),

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

@ -182,7 +182,9 @@ TEST_P(MslIntrinsicTest, Emit) {
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);

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

@ -168,10 +168,10 @@ type::Matrix* GetNestedMatrixType(type::Type* type) {
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 @@ 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<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 @@ 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<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 @@ 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<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 @@ 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<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 @@ 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<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 @@ 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<type::MultisampledTexture>() ||
texture_type->Is<type::StorageTexture>()) {
@ -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<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 @@ 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<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 @@ 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;
}

19
src/writer/spirv/builder.h
View File

@ -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