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

sem: Split sem::Variable into global, local and parameter 

Each of these may contain information specific to their kind.

Change-Id: Ic8ac808088132b7bc2e43da6ce46a06571e0fed5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/59200
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ryan Harrison <rharrison@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
This commit is contained in:
Ben Clayton 2021-07-22 13:24:59 +00:00
parent 66b979d7fb
commit 0f2d95dea3
36 changed files with 435 additions and 277 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
src/ast/identifier_expression.h
View File

@ -16,7 +16,6 @@
#define SRC_AST_IDENTIFIER_EXPRESSION_H_
#include "src/ast/expression.h"
#include "src/sem/intrinsic.h"
namespace tint {
namespace ast {

15
src/inspector/inspector.cc
View File

@ -174,7 +174,8 @@ std::vector<EntryPoint> Inspector::GetEntryPoints() {
auto name = program_->Symbols().NameFor(decl->symbol());
if (var->IsPipelineConstant()) {
auto* global = var->As<sem::GlobalVariable>();
if (global && global->IsPipelineConstant()) {
OverridableConstant overridable_constant;
overridable_constant.name = name;
entry_point.overridable_constants.push_back(overridable_constant);
@ -203,8 +204,8 @@ std::string Inspector::GetRemappedNameForEntryPoint(
std::map<uint32_t, Scalar> Inspector::GetConstantIDs() {
std::map<uint32_t, Scalar> result;
for (auto* var : program_->AST().GlobalVariables()) {
auto* sem_var = program_->Sem().Get(var);
if (!sem_var->IsPipelineConstant()) {
auto* global = program_->Sem().Get<sem::GlobalVariable>(var);
if (!global || !global->IsPipelineConstant()) {
continue;
}
@ -212,7 +213,7 @@ std::map<uint32_t, Scalar> Inspector::GetConstantIDs() {
// WGSL, so the resolver should catch it. Thus here the inspector just
// assumes all definitions of the constant id are the same, so only needs
// to find the first reference to constant id.
uint32_t constant_id = sem_var->ConstantId();
uint32_t constant_id = global->ConstantId();
if (result.find(constant_id) != result.end()) {
continue;
}
@ -274,10 +275,10 @@ std::map<uint32_t, Scalar> Inspector::GetConstantIDs() {
std::map<std::string, uint32_t> Inspector::GetConstantNameToIdMap() {
std::map<std::string, uint32_t> result;
for (auto* var : program_->AST().GlobalVariables()) {
auto* sem_var = program_->Sem().Get(var);
if (sem_var->IsPipelineConstant()) {
auto* global = program_->Sem().Get<sem::GlobalVariable>(var);
if (global && global->IsPipelineConstant()) {
auto name = program_->Symbols().NameFor(var->symbol());
result[name] = sem_var->ConstantId();
result[name] = global->ConstantId();
}
}
return result;

15
src/inspector/inspector_test.cc
View File

@ -928,16 +928,19 @@ TEST_F(InspectorGetConstantNameToIdMapTest, WithAndWithoutIds) {
EXPECT_EQ(result["v300"], 300u);
ASSERT_TRUE(result.count("a"));
ASSERT_TRUE(program_->Sem().Get(a));
EXPECT_EQ(result["a"], program_->Sem().Get(a)->ConstantId());
ASSERT_TRUE(program_->Sem().Get<sem::GlobalVariable>(a));
EXPECT_EQ(result["a"],
program_->Sem().Get<sem::GlobalVariable>(a)->ConstantId());
ASSERT_TRUE(result.count("b"));
ASSERT_TRUE(program_->Sem().Get(b));
EXPECT_EQ(result["b"], program_->Sem().Get(b)->ConstantId());
ASSERT_TRUE(program_->Sem().Get<sem::GlobalVariable>(b));
EXPECT_EQ(result["b"],
program_->Sem().Get<sem::GlobalVariable>(b)->ConstantId());
ASSERT_TRUE(result.count("c"));
ASSERT_TRUE(program_->Sem().Get(c));
EXPECT_EQ(result["c"], program_->Sem().Get(c)->ConstantId());
ASSERT_TRUE(program_->Sem().Get<sem::GlobalVariable>(c));
EXPECT_EQ(result["c"],
program_->Sem().Get<sem::GlobalVariable>(c)->ConstantId());
}
TEST_F(InspectorGetStorageSizeTest, Empty) {

80
src/intrinsic_table.cc
View File

@ -701,6 +701,55 @@ struct IntrinsicInfo {
#include "intrinsic_table.inl"
/// IntrinsicPrototype describes a fully matched intrinsic function, which is
/// used as a lookup for building unique sem::Intrinsic instances.
struct IntrinsicPrototype {
/// Parameter describes a single parameter
struct Parameter {
/// Parameter type
sem::Type* const type;
/// Parameter usage
ParameterUsage const usage = ParameterUsage::kNone;
};
/// Hasher provides a hash function for the IntrinsicPrototype
struct Hasher {
/// @param i the IntrinsicPrototype to create a hash for
/// @return the hash value
inline std::size_t operator()(const IntrinsicPrototype& i) const {
size_t hash = utils::Hash(i.parameters.size());
for (auto& p : i.parameters) {
utils::HashCombine(&hash, p.type, p.usage);
}
return utils::Hash(hash, i.type, i.return_type, i.supported_stages,
i.is_deprecated);
}
};
sem::IntrinsicType type = sem::IntrinsicType::kNone;
std::vector<Parameter> parameters;
sem::Type const* return_type = nullptr;
PipelineStageSet supported_stages;
bool is_deprecated = false;
};
/// Equality operator for IntrinsicPrototype
bool operator==(const IntrinsicPrototype& a, const IntrinsicPrototype& b) {
if (a.type != b.type || a.supported_stages != b.supported_stages ||
a.return_type != b.return_type || a.is_deprecated != b.is_deprecated ||
a.parameters.size() != b.parameters.size()) {
return false;
}
for (size_t i = 0; i < a.parameters.size(); i++) {
auto& pa = a.parameters[i];
auto& pb = b.parameters[i];
if (pa.type != pb.type || pa.usage != pb.usage) {
return false;
}
}
return true;
}
/// Impl is the private implementation of the IntrinsicTable interface.
class Impl : public IntrinsicTable {
public:
@ -726,7 +775,10 @@ class Impl : public IntrinsicTable {
ProgramBuilder& builder;
Matchers matchers;
std::unordered_map<sem::Intrinsic, sem::Intrinsic*> intrinsics;
std::unordered_map<IntrinsicPrototype,
sem::Intrinsic*,
IntrinsicPrototype::Hasher>
intrinsics;
};
/// @return a string representing a call to an intrinsic with the given argument
@ -833,7 +885,7 @@ const sem::Intrinsic* Impl::Match(sem::IntrinsicType intrinsic_type,
ClosedState closed(builder);
sem::ParameterList parameters;
std::vector<IntrinsicPrototype::Parameter> parameters;
auto num_params = std::min(num_parameters, num_arguments);
for (uint32_t p = 0; p < num_params; p++) {
@ -841,8 +893,8 @@ const sem::Intrinsic* Impl::Match(sem::IntrinsicType intrinsic_type,
auto* indices = parameter.matcher_indices;
auto* type = Match(closed, overload, indices).Type(args[p]->UnwrapRef());
if (type) {
parameters.emplace_back(
sem::Parameter{const_cast<sem::Type*>(type), parameter.usage});
parameters.emplace_back(IntrinsicPrototype::Parameter{
const_cast<sem::Type*>(type), parameter.usage});
match_score += kScorePerMatchedParam;
} else {
overload_matched = false;
@ -899,13 +951,25 @@ const sem::Intrinsic* Impl::Match(sem::IntrinsicType intrinsic_type,
return_type = builder.create<sem::Void>();
}
sem::Intrinsic intrinsic(intrinsic_type, const_cast<sem::Type*>(return_type),
std::move(parameters), overload.supported_stages,
overload.is_deprecated);
IntrinsicPrototype intrinsic;
intrinsic.type = intrinsic_type;
intrinsic.return_type = return_type;
intrinsic.parameters = std::move(parameters);
intrinsic.supported_stages = overload.supported_stages;
intrinsic.is_deprecated = overload.is_deprecated;
// De-duplicate intrinsics that are identical.
return utils::GetOrCreate(intrinsics, intrinsic, [&] {
return builder.create<sem::Intrinsic>(intrinsic);
sem::ParameterList params;
params.reserve(intrinsic.parameters.size());
for (auto& p : intrinsic.parameters) {
params.emplace_back(builder.create<sem::Parameter>(
nullptr, p.type, ast::StorageClass::kNone, ast::Access::kUndefined,
p.usage));
}
return builder.create<sem::Intrinsic>(
intrinsic.type, const_cast<sem::Type*>(intrinsic.return_type),
std::move(params), intrinsic.supported_stages, intrinsic.is_deprecated);
});
}

143
src/intrinsic_table_test.cc
View File

@ -45,7 +45,8 @@ TEST_F(IntrinsicTableTest, MatchF32) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kCos);
EXPECT_THAT(result->ReturnType(), f32);
EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{f32}));
ASSERT_THAT(result->Parameters().size(), 1);
EXPECT_EQ(result->Parameters()[0]->Type(), f32);
}
TEST_F(IntrinsicTableTest, MismatchF32) {
@ -65,7 +66,8 @@ TEST_F(IntrinsicTableTest, MatchU32) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kUnpack2x16float);
EXPECT_THAT(result->ReturnType(), vec2_f32);
EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{u32}));
ASSERT_EQ(result->Parameters().size(), 1u);
EXPECT_EQ(result->Parameters()[0]->Type(), u32);
}
TEST_F(IntrinsicTableTest, MismatchU32) {
@ -87,10 +89,13 @@ TEST_F(IntrinsicTableTest, MatchI32) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
EXPECT_THAT(result->ReturnType(), vec4_f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{tex, ParameterUsage::kTexture},
Parameter{i32, ParameterUsage::kCoords},
Parameter{i32, ParameterUsage::kLevel}));
ASSERT_EQ(result->Parameters().size(), 3u);
EXPECT_EQ(result->Parameters()[0]->Type(), tex);
EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
EXPECT_EQ(result->Parameters()[1]->Type(), i32);
EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
EXPECT_EQ(result->Parameters()[2]->Type(), i32);
EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kLevel);
}
TEST_F(IntrinsicTableTest, MismatchI32) {
@ -109,7 +114,8 @@ TEST_F(IntrinsicTableTest, MatchIU32AsI32) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kCountOneBits);
EXPECT_THAT(result->ReturnType(), i32);
EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{i32}));
ASSERT_EQ(result->Parameters().size(), 1u);
EXPECT_EQ(result->Parameters()[0]->Type(), i32);
}
TEST_F(IntrinsicTableTest, MatchIU32AsU32) {
@ -119,7 +125,8 @@ TEST_F(IntrinsicTableTest, MatchIU32AsU32) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kCountOneBits);
EXPECT_THAT(result->ReturnType(), u32);
EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{u32}));
ASSERT_EQ(result->Parameters().size(), 1u);
EXPECT_EQ(result->Parameters()[0]->Type(), u32);
}
TEST_F(IntrinsicTableTest, MismatchIU32) {
@ -137,8 +144,10 @@ TEST_F(IntrinsicTableTest, MatchFIU32AsI32) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kClamp);
EXPECT_THAT(result->ReturnType(), i32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{i32}, Parameter{i32}, Parameter{i32}));
ASSERT_EQ(result->Parameters().size(), 3u);
EXPECT_EQ(result->Parameters()[0]->Type(), i32);
EXPECT_EQ(result->Parameters()[1]->Type(), i32);
EXPECT_EQ(result->Parameters()[2]->Type(), i32);
}
TEST_F(IntrinsicTableTest, MatchFIU32AsU32) {
@ -149,8 +158,10 @@ TEST_F(IntrinsicTableTest, MatchFIU32AsU32) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kClamp);
EXPECT_THAT(result->ReturnType(), u32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{u32}, Parameter{u32}, Parameter{u32}));
ASSERT_EQ(result->Parameters().size(), 3u);
EXPECT_EQ(result->Parameters()[0]->Type(), u32);
EXPECT_EQ(result->Parameters()[1]->Type(), u32);
EXPECT_EQ(result->Parameters()[2]->Type(), u32);
}
TEST_F(IntrinsicTableTest, MatchFIU32AsF32) {
@ -161,8 +172,10 @@ TEST_F(IntrinsicTableTest, MatchFIU32AsF32) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kClamp);
EXPECT_THAT(result->ReturnType(), f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{f32}, Parameter{f32}, Parameter{f32}));
ASSERT_EQ(result->Parameters().size(), 3u);
EXPECT_EQ(result->Parameters()[0]->Type(), f32);
EXPECT_EQ(result->Parameters()[1]->Type(), f32);
EXPECT_EQ(result->Parameters()[2]->Type(), f32);
}
TEST_F(IntrinsicTableTest, MismatchFIU32) {
@ -182,8 +195,10 @@ TEST_F(IntrinsicTableTest, MatchBool) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kSelect);
EXPECT_THAT(result->ReturnType(), f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{f32}, Parameter{f32}, Parameter{bool_}));
ASSERT_EQ(result->Parameters().size(), 3u);
EXPECT_EQ(result->Parameters()[0]->Type(), f32);
EXPECT_EQ(result->Parameters()[1]->Type(), f32);
EXPECT_EQ(result->Parameters()[2]->Type(), bool_);
}
TEST_F(IntrinsicTableTest, MismatchBool) {
@ -203,8 +218,9 @@ TEST_F(IntrinsicTableTest, MatchPointer) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kModf);
EXPECT_THAT(result->ReturnType(), f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{f32}, Parameter{ptr}));
ASSERT_EQ(result->Parameters().size(), 2u);
EXPECT_EQ(result->Parameters()[0]->Type(), f32);
EXPECT_EQ(result->Parameters()[1]->Type(), ptr);
}
TEST_F(IntrinsicTableTest, MismatchPointer) {
@ -225,7 +241,7 @@ TEST_F(IntrinsicTableTest, MatchArray) {
EXPECT_THAT(result->Type(), IntrinsicType::kArrayLength);
EXPECT_TRUE(result->ReturnType()->Is<sem::U32>());
ASSERT_EQ(result->Parameters().size(), 1u);
auto* param_type = result->Parameters()[0].type;
auto* param_type = result->Parameters()[0]->Type();
ASSERT_TRUE(param_type->Is<sem::Pointer>());
EXPECT_TRUE(param_type->As<sem::Pointer>()->StoreType()->Is<sem::Array>());
}
@ -249,10 +265,13 @@ TEST_F(IntrinsicTableTest, MatchSampler) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kTextureSample);
EXPECT_THAT(result->ReturnType(), vec4_f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{tex, ParameterUsage::kTexture},
Parameter{sampler, ParameterUsage::kSampler},
Parameter{vec2_f32, ParameterUsage::kCoords}));
ASSERT_EQ(result->Parameters().size(), 3u);
EXPECT_EQ(result->Parameters()[0]->Type(), tex);
EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
EXPECT_EQ(result->Parameters()[1]->Type(), sampler);
EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kSampler);
EXPECT_EQ(result->Parameters()[2]->Type(), vec2_f32);
EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kCoords);
}
TEST_F(IntrinsicTableTest, MismatchSampler) {
@ -277,10 +296,13 @@ TEST_F(IntrinsicTableTest, MatchSampledTexture) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
EXPECT_THAT(result->ReturnType(), vec4_f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{tex, ParameterUsage::kTexture},
Parameter{vec2_i32, ParameterUsage::kCoords},
Parameter{i32, ParameterUsage::kLevel}));
ASSERT_EQ(result->Parameters().size(), 3u);
EXPECT_EQ(result->Parameters()[0]->Type(), tex);
EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
EXPECT_EQ(result->Parameters()[2]->Type(), i32);
EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kLevel);
}
TEST_F(IntrinsicTableTest, MatchMultisampledTexture) {
@ -295,10 +317,13 @@ TEST_F(IntrinsicTableTest, MatchMultisampledTexture) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
EXPECT_THAT(result->ReturnType(), vec4_f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{tex, ParameterUsage::kTexture},
Parameter{vec2_i32, ParameterUsage::kCoords},
Parameter{i32, ParameterUsage::kSampleIndex}));
ASSERT_EQ(result->Parameters().size(), 3u);
EXPECT_EQ(result->Parameters()[0]->Type(), tex);
EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
EXPECT_EQ(result->Parameters()[2]->Type(), i32);
EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kSampleIndex);
}
TEST_F(IntrinsicTableTest, MatchDepthTexture) {
@ -312,10 +337,13 @@ TEST_F(IntrinsicTableTest, MatchDepthTexture) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
EXPECT_THAT(result->ReturnType(), f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{tex, ParameterUsage::kTexture},
Parameter{vec2_i32, ParameterUsage::kCoords},
Parameter{i32, ParameterUsage::kLevel}));
ASSERT_EQ(result->Parameters().size(), 3u);
EXPECT_EQ(result->Parameters()[0]->Type(), tex);
EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
EXPECT_EQ(result->Parameters()[2]->Type(), i32);
EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kLevel);
}
TEST_F(IntrinsicTableTest, MatchExternalTexture) {
@ -330,9 +358,11 @@ TEST_F(IntrinsicTableTest, MatchExternalTexture) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
EXPECT_THAT(result->ReturnType(), vec4_f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{tex, ParameterUsage::kTexture},
Parameter{vec2_i32, ParameterUsage::kCoords}));
ASSERT_EQ(result->Parameters().size(), 2u);
EXPECT_EQ(result->Parameters()[0]->Type(), tex);
EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
}
TEST_F(IntrinsicTableTest, MatchROStorageTexture) {
@ -352,9 +382,11 @@ TEST_F(IntrinsicTableTest, MatchROStorageTexture) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kTextureLoad);
EXPECT_THAT(result->ReturnType(), vec4_f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{tex, ParameterUsage::kTexture},
Parameter{vec2_i32, ParameterUsage::kCoords}));
ASSERT_EQ(result->Parameters().size(), 2u);
EXPECT_EQ(result->Parameters()[0]->Type(), tex);
EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
}
TEST_F(IntrinsicTableTest, MatchWOStorageTexture) {
@ -374,10 +406,13 @@ TEST_F(IntrinsicTableTest, MatchWOStorageTexture) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kTextureStore);
EXPECT_TRUE(result->ReturnType()->Is<sem::Void>());
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{tex, ParameterUsage::kTexture},
Parameter{vec2_i32, ParameterUsage::kCoords},
Parameter{vec4_f32, ParameterUsage::kValue}));
ASSERT_EQ(result->Parameters().size(), 3u);
EXPECT_EQ(result->Parameters()[0]->Type(), tex);
EXPECT_EQ(result->Parameters()[0]->Usage(), ParameterUsage::kTexture);
EXPECT_EQ(result->Parameters()[1]->Type(), vec2_i32);
EXPECT_EQ(result->Parameters()[1]->Usage(), ParameterUsage::kCoords);
EXPECT_EQ(result->Parameters()[2]->Type(), vec4_f32);
EXPECT_EQ(result->Parameters()[2]->Usage(), ParameterUsage::kValue);
}
TEST_F(IntrinsicTableTest, MismatchTexture) {
@ -401,7 +436,8 @@ TEST_F(IntrinsicTableTest, ImplicitLoadOnReference) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kCos);
EXPECT_THAT(result->ReturnType(), f32);
EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{f32}));
ASSERT_EQ(result->Parameters().size(), 1u);
EXPECT_EQ(result->Parameters()[0]->Type(), f32);
}
TEST_F(IntrinsicTableTest, MatchOpenType) {
@ -412,8 +448,9 @@ TEST_F(IntrinsicTableTest, MatchOpenType) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kClamp);
EXPECT_THAT(result->ReturnType(), f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{f32}, Parameter{f32}, Parameter{f32}));
EXPECT_EQ(result->Parameters()[0]->Type(), f32);
EXPECT_EQ(result->Parameters()[1]->Type(), f32);
EXPECT_EQ(result->Parameters()[2]->Type(), f32);
}
TEST_F(IntrinsicTableTest, MismatchOpenType) {
@ -434,9 +471,10 @@ TEST_F(IntrinsicTableTest, MatchOpenSizeVector) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kClamp);
EXPECT_THAT(result->ReturnType(), vec2_f32);
EXPECT_THAT(result->Parameters(),
ElementsAre(Parameter{vec2_f32}, Parameter{vec2_f32},
Parameter{vec2_f32}));
ASSERT_EQ(result->Parameters().size(), 3u);
EXPECT_EQ(result->Parameters()[0]->Type(), vec2_f32);
EXPECT_EQ(result->Parameters()[1]->Type(), vec2_f32);
EXPECT_EQ(result->Parameters()[2]->Type(), vec2_f32);
}
TEST_F(IntrinsicTableTest, MismatchOpenSizeVector) {
@ -459,7 +497,8 @@ TEST_F(IntrinsicTableTest, MatchOpenSizeMatrix) {
ASSERT_EQ(Diagnostics().str(), "");
EXPECT_THAT(result->Type(), IntrinsicType::kDeterminant);
EXPECT_THAT(result->ReturnType(), f32);
EXPECT_THAT(result->Parameters(), ElementsAre(Parameter{mat3_f32}));
ASSERT_EQ(result->Parameters().size(), 1u);
EXPECT_EQ(result->Parameters()[0]->Type(), mat3_f32);
}
TEST_F(IntrinsicTableTest, MismatchOpenSizeMatrix) {

1
src/program_builder.h
View File

@ -50,6 +50,7 @@
#include "src/ast/pointer.h"
#include "src/ast/return_statement.h"
#include "src/ast/sampled_texture.h"
#include "src/ast/sampler.h"
#include "src/ast/scalar_constructor_expression.h"
#include "src/ast/sint_literal.h"
#include "src/ast/stage_decoration.h"

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

@ -34,6 +34,7 @@
#include "src/ast/unary_op_expression.h"
#include "src/ast/variable_decl_statement.h"
#include "src/sem/depth_texture_type.h"
#include "src/sem/intrinsic_type.h"
#include "src/sem/sampled_texture_type.h"
// Terms:
@ -4959,9 +4960,7 @@ bool FunctionEmitter::EmitControlBarrier(
TypedExpression FunctionEmitter::MakeIntrinsicCall(
const spvtools::opt::Instruction& inst) {
const auto intrinsic = GetIntrinsic(inst.opcode());
std::ostringstream ss;
ss << intrinsic;
auto name = ss.str();
auto* name = sem::str(intrinsic);
auto* ident = create<ast::IdentifierExpression>(
Source{}, builder_.Symbols().Register(name));

4
src/resolver/intrinsic_test.cc
View File

@ -1706,11 +1706,11 @@ std::string to_str(const std::string& function,
std::stringstream out;
out << function << "(";
bool first = true;
for (auto& param : params) {
for (auto* param : params) {
if (!first) {
out << ", ";
}
out << sem::str(param.usage);
out << sem::str(param->Usage());
first = false;
}
out << ")";

8
src/resolver/pipeline_overridable_constant_test.cc
View File

@ -30,7 +30,7 @@ TEST_F(ResolverPipelineOverridableConstantTest, NonOverridable) {
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* sem_a = Sem().Get(a);
auto* sem_a = Sem().Get<sem::GlobalVariable>(a);
ASSERT_NE(sem_a, nullptr);
EXPECT_EQ(sem_a->Declaration(), a);
EXPECT_FALSE(sem_a->IsPipelineConstant());
@ -41,7 +41,7 @@ TEST_F(ResolverPipelineOverridableConstantTest, WithId) {
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* sem_a = Sem().Get(a);
auto* sem_a = Sem().Get<sem::GlobalVariable>(a);
ASSERT_NE(sem_a, nullptr);
EXPECT_EQ(sem_a->Declaration(), a);
EXPECT_TRUE(sem_a->IsPipelineConstant());
@ -53,7 +53,7 @@ TEST_F(ResolverPipelineOverridableConstantTest, WithoutId) {
EXPECT_TRUE(r()->Resolve()) << r()->error();
auto* sem_a = Sem().Get(a);
auto* sem_a = Sem().Get<sem::GlobalVariable>(a);
ASSERT_NE(sem_a, nullptr);
EXPECT_EQ(sem_a->Declaration(), a);
EXPECT_TRUE(sem_a->IsPipelineConstant());
@ -79,7 +79,7 @@ TEST_F(ResolverPipelineOverridableConstantTest, WithAndWithoutIds) {
std::vector<uint16_t> constant_ids;
for (auto* var : variables) {
auto* sem = Sem().Get(var);
auto* sem = Sem().Get<sem::GlobalVariable>(var);
ASSERT_NE(sem, nullptr);
constant_ids.push_back(static_cast<uint16_t>(sem->ConstantId()));
}

29
src/resolver/resolver.cc
View File

@ -3693,11 +3693,24 @@ void Resolver::CreateSemanticNodes() const {
next_constant_id = constant_id + 1;
}
sem_var = builder_->create<sem::Variable>(var, info->type, constant_id);
} else {
sem_var =
builder_->create<sem::Variable>(var, info->type, info->storage_class,
info->access, info->binding_point);
builder_->create<sem::GlobalVariable>(var, info->type, constant_id);
} else {
switch (info->kind) {
case VariableKind::kGlobal:
sem_var = builder_->create<sem::GlobalVariable>(
var, info->type, info->storage_class, info->access,
info->binding_point);
break;
case VariableKind::kLocal:
sem_var = builder_->create<sem::LocalVariable>(
var, info->type, info->storage_class, info->access);
break;
case VariableKind::kParameter:
sem_var = builder_->create<sem::Parameter>(
var, info->type, info->storage_class, info->access);
break;
}
}
std::vector<const sem::VariableUser*> users;
@ -3739,9 +3752,15 @@ void Resolver::CreateSemanticNodes() const {
auto* func = it.first;
auto* info = it.second;
sem::ParameterList parameters;
parameters.reserve(info->parameters.size());
for (auto* p : info->parameters) {
parameters.emplace_back(sem.Get<sem::Parameter>(p->declaration));
}
auto* sem_func = builder_->create<sem::Function>(
info->declaration, const_cast<sem::Type*>(info->return_type),
remap_vars(info->parameters), remap_vars(info->referenced_module_vars),
parameters, remap_vars(info->referenced_module_vars),
remap_vars(info->local_referenced_module_vars), info->return_statements,
info->callsites, ancestor_entry_points[func->symbol()],
info->workgroup_size);

6
src/resolver/resolver_test.cc
View File

@ -2018,8 +2018,10 @@ TEST_F(ResolverTest, BindingPoint_SetForResources) {
EXPECT_TRUE(r()->Resolve()) << r()->error();
EXPECT_EQ(Sem().Get(s1)->BindingPoint(), (sem::BindingPoint{1u, 2u}));
EXPECT_EQ(Sem().Get(s2)->BindingPoint(), (sem::BindingPoint{3u, 4u}));
EXPECT_EQ(Sem().Get<sem::GlobalVariable>(s1)->BindingPoint(),
(sem::BindingPoint{1u, 2u}));
EXPECT_EQ(Sem().Get<sem::GlobalVariable>(s2)->BindingPoint(),
(sem::BindingPoint{3u, 4u}));
}
TEST_F(ResolverTest, Function_EntryPoints_StageDecoration) {

8
src/sem/call_target.cc
View File

@ -32,18 +32,12 @@ CallTarget::~CallTarget() = default;
int IndexOf(const ParameterList& parameters, ParameterUsage usage) {
for (size_t i = 0; i < parameters.size(); i++) {
if (parameters[i].usage == usage) {
if (parameters[i]->Usage() == usage) {
return static_cast<int>(i);
}
}
return -1;
}
std::ostream& operator<<(std::ostream& out, Parameter parameter) {
out << "[type: " << parameter.type->FriendlyName(SymbolTable{ProgramID{}})
<< ", usage: " << str(parameter.usage) << "]";
return out;
}
} // namespace sem
} // namespace tint

41
src/sem/call_target.h
View File

@ -18,39 +18,15 @@
#include <vector>
#include "src/sem/node.h"
#include "src/sem/parameter_usage.h"
#include "src/sem/sampler_type.h"
#include "src/sem/variable.h"
#include "src/utils/hash.h"
namespace tint {
namespace sem {
// Forward declarations
class Type;
/// Parameter describes a single parameter of a call target
struct Parameter {
/// Parameter type
sem::Type* const type;
/// Parameter usage
ParameterUsage const usage = ParameterUsage::kNone;
};
std::ostream& operator<<(std::ostream& out, Parameter parameter);
/// Equality operator for Parameters
static inline bool operator==(const Parameter& a, const Parameter& b) {
return a.type == b.type && a.usage == b.usage;
}
/// Inequality operator for Parameters
static inline bool operator!=(const Parameter& a, const Parameter& b) {
return !(a == b);
}
/// ParameterList is a list of Parameter
using ParameterList = 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
@ -85,19 +61,4 @@ class CallTarget : public Castable<CallTarget, Node> {
} // namespace sem
} // namespace tint
namespace std {
/// Custom std::hash specialization for tint::sem::Parameter
template <>
class hash<tint::sem::Parameter> {
public:
/// @param p the tint::sem::Parameter to create a hash for
/// @return the hash value
inline std::size_t operator()(const tint::sem::Parameter& p) const {
return tint::utils::Hash(p.type, p.usage);
}
};
} // namespace std
#endif // SRC_SEM_CALL_TARGET_H_

18
src/sem/function.cc
View File

@ -27,31 +27,17 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::Function);
namespace tint {
namespace sem {
namespace {
ParameterList GetParameters(const std::vector<const Variable*>& params) {
ParameterList parameters;
parameters.reserve(params.size());
for (auto* param : params) {
parameters.emplace_back(Parameter{param->Type(), ParameterUsage::kNone});
}
return parameters;
}
} // namespace
Function::Function(ast::Function* declaration,
Type* return_type,
std::vector<const Variable*> parameters,
ParameterList parameters,
std::vector<const Variable*> referenced_module_vars,
std::vector<const Variable*> local_referenced_module_vars,
std::vector<const ast::ReturnStatement*> return_statements,
std::vector<const ast::CallExpression*> callsites,
std::vector<Symbol> ancestor_entry_points,
std::array<WorkgroupDimension, 3> workgroup_size)
: Base(return_type, GetParameters(parameters)),
: Base(return_type, std::move(parameters)),
declaration_(declaration),
parameters_(std::move(parameters)),
referenced_module_vars_(std::move(referenced_module_vars)),
local_referenced_module_vars_(std::move(local_referenced_module_vars)),
return_statements_(std::move(return_statements)),

6
src/sem/function.h
View File

@ -68,7 +68,7 @@ class Function : public Castable<Function, CallTarget> {
/// @param workgroup_size the workgroup size
Function(ast::Function* declaration,
Type* return_type,
std::vector<const Variable*> parameters,
ParameterList parameters,
std::vector<const Variable*> referenced_module_vars,
std::vector<const Variable*> local_referenced_module_vars,
std::vector<const ast::ReturnStatement*> return_statements,
@ -82,9 +82,6 @@ class Function : public Castable<Function, CallTarget> {
/// @returns the ast::Function declaration
ast::Function* Declaration() const { return declaration_; }
/// @return the parameters to the function
const std::vector<const Variable*> Parameters() const { return parameters_; }
/// Note: If this function calls other functions, the return will also include
/// all of the referenced variables from the callees.
/// @returns the referenced module variables
@ -178,7 +175,6 @@ class Function : public Castable<Function, CallTarget> {
bool multisampled) const;
ast::Function* const declaration_;
std::vector<const Variable*> const parameters_;
std::vector<const Variable*> const referenced_module_vars_;
std::vector<const Variable*> const local_referenced_module_vars_;
std::vector<const ast::ReturnStatement*> const return_statements_;

29
src/sem/intrinsic.cc
View File

@ -19,11 +19,6 @@ TINT_INSTANTIATE_TYPEINFO(tint::sem::Intrinsic);
namespace tint {
namespace sem {
std::ostream& operator<<(std::ostream& out, IntrinsicType i) {
out << str(i);
return out;
}
const char* Intrinsic::str() const {
return sem::str(type_);
}
@ -103,7 +98,7 @@ bool IsAtomicIntrinsic(IntrinsicType i) {
Intrinsic::Intrinsic(IntrinsicType type,
sem::Type* return_type,
const ParameterList& parameters,
ParameterList parameters,
PipelineStageSet supported_stages,
bool is_deprecated)
: Base(return_type, parameters),
@ -111,8 +106,6 @@ Intrinsic::Intrinsic(IntrinsicType type,
supported_stages_(supported_stages),
is_deprecated_(is_deprecated) {}
Intrinsic::Intrinsic(const Intrinsic&) = default;
Intrinsic::~Intrinsic() = default;
bool Intrinsic::IsCoarseDerivative() const {
@ -155,25 +148,5 @@ bool Intrinsic::IsAtomic() const {
return IsAtomicIntrinsic(type_);
}
bool operator==(const Intrinsic& a, const Intrinsic& b) {
static_assert(sizeof(Intrinsic(IntrinsicType::kNone, nullptr, ParameterList{},
PipelineStageSet{}, false)) > 0,
"don't forget to update the comparison below if you change the "
"constructor of Intrinsic!");
if (a.Type() != b.Type() || a.SupportedStages() != b.SupportedStages() ||
a.ReturnType() != b.ReturnType() ||
a.IsDeprecated() != b.IsDeprecated() ||
a.Parameters().size() != b.Parameters().size()) {
return false;
}
for (size_t i = 0; i < a.Parameters().size(); i++) {
if (a.Parameters()[i] != b.Parameters()[i]) {
return false;
}
}
return true;
}
} // namespace sem
} // namespace tint

17
src/sem/intrinsic.h
View File

@ -88,13 +88,10 @@ class Intrinsic : public Castable<Intrinsic, CallTarget> {
/// deprecated
Intrinsic(IntrinsicType type,
sem::Type* return_type,
const ParameterList& parameters,
ParameterList parameters,
PipelineStageSet supported_stages,
bool is_deprecated);
/// Copy constructor
Intrinsic(const Intrinsic&);
/// Destructor
~Intrinsic() override;
@ -147,18 +144,6 @@ class Intrinsic : public Castable<Intrinsic, CallTarget> {
bool const is_deprecated_;
};
/// 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);
/// Equality operator for Intrinsics
bool operator==(const Intrinsic& a, const Intrinsic& b);
/// Inequality operator for Intrinsics
static inline bool operator!=(const Intrinsic& a, const Intrinsic& b) {
return !(a == b);
}
} // namespace sem
} // namespace tint

7
src/sem/intrinsic_type.cc
View File

@ -24,6 +24,8 @@
#include "src/sem/intrinsic_type.h"
#include <sstream>
namespace tint {
namespace sem {
@ -529,5 +531,10 @@ const char* str(IntrinsicType i) {
return "<unknown>";
}
std::ostream& operator<<(std::ostream& out, IntrinsicType i) {
out << str(i);
return out;
}
} // namespace sem
} // namespace tint

7
src/sem/intrinsic_type.cc.tmpl
View File

@ -10,6 +10,8 @@ See:
#include "src/sem/intrinsic_type.h"
#include <sstream>
namespace tint {
namespace sem {
@ -34,5 +36,10 @@ const char* str(IntrinsicType i) {
return "<unknown>";
}
std::ostream& operator<<(std::ostream& out, IntrinsicType i) {
out << str(i);
return out;
}
} // namespace sem
} // namespace tint

5
src/sem/intrinsic_type.h
View File

@ -26,6 +26,7 @@
#define SRC_SEM_INTRINSIC_TYPE_H_
#include <string>
#include <sstream>
namespace tint {
namespace sem {
@ -143,6 +144,10 @@ IntrinsicType ParseIntrinsicType(const std::string& name);
/// case, matches the name in the WGSL spec.
const char* str(IntrinsicType i);
/// 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 sem
} // namespace tint

5
src/sem/intrinsic_type.h.tmpl
View File

@ -12,6 +12,7 @@ See:
#define SRC_SEM_INTRINSIC_TYPE_H_
#include <string>
#include <sstream>
namespace tint {
namespace sem {
@ -34,6 +35,10 @@ IntrinsicType ParseIntrinsicType(const std::string& name);
/// case, matches the name in the WGSL spec.
const char* str(IntrinsicType i);
/// 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 sem
} // namespace tint

1
src/sem/sampler_type_test.cc
View File

@ -12,6 +12,7 @@
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/sem/sampler_type.h"
#include "src/sem/test_helper.h"
#include "src/sem/texture_type.h"

46
src/sem/variable.cc
View File

@ -20,6 +20,9 @@
#include "src/ast/variable.h"
TINT_INSTANTIATE_TYPEINFO(tint::sem::Variable);
TINT_INSTANTIATE_TYPEINFO(tint::sem::GlobalVariable);
TINT_INSTANTIATE_TYPEINFO(tint::sem::LocalVariable);
TINT_INSTANTIATE_TYPEINFO(tint::sem::Parameter);
TINT_INSTANTIATE_TYPEINFO(tint::sem::VariableUser);
namespace tint {
@ -28,26 +31,51 @@ namespace sem {
Variable::Variable(const ast::Variable* declaration,
const sem::Type* type,
ast::StorageClass storage_class,
ast::Access access,
sem::BindingPoint binding_point)
ast::Access access)
: declaration_(declaration),
type_(type),
storage_class_(storage_class),
access_(access),
access_(access) {}
Variable::~Variable() = default;
LocalVariable::LocalVariable(const ast::Variable* declaration,
const sem::Type* type,
ast::StorageClass storage_class,
ast::Access access)
: Base(declaration, type, storage_class, access) {}
LocalVariable::~LocalVariable() = default;
GlobalVariable::GlobalVariable(const ast::Variable* declaration,
const sem::Type* type,
ast::StorageClass storage_class,
ast::Access access,
sem::BindingPoint binding_point)
: Base(declaration, type, storage_class, access),
binding_point_(binding_point),
is_pipeline_constant_(false) {}
Variable::Variable(const ast::Variable* declaration,
GlobalVariable::GlobalVariable(const ast::Variable* declaration,
const sem::Type* type,
uint16_t constant_id)
: declaration_(declaration),
type_(type),
storage_class_(ast::StorageClass::kNone),
access_(ast::Access::kReadWrite),
: Base(declaration,
type,
ast::StorageClass::kNone,
ast::Access::kReadWrite),
is_pipeline_constant_(true),
constant_id_(constant_id) {}
Variable::~Variable() = default;
GlobalVariable::~GlobalVariable() = default;
Parameter::Parameter(const ast::Variable* declaration,
const sem::Type* type,
ast::StorageClass storage_class,
ast::Access access,
const ParameterUsage usage /* = ParameterUsage::kNone */)
: Base(declaration, type, storage_class, access), usage_(usage) {}
Parameter::~Parameter() = default;
VariableUser::VariableUser(ast::IdentifierExpression* declaration,
const sem::Type* type,

119
src/sem/variable.h
View File

@ -21,6 +21,7 @@
#include "src/ast/storage_class.h"
#include "src/sem/binding_point.h"
#include "src/sem/expression.h"
#include "src/sem/parameter_usage.h"
namespace tint {
@ -36,28 +37,19 @@ namespace sem {
class Type;
class VariableUser;
/// Variable holds the semantic information for variables.
/// Variable is the base class for local variables, global variables and
/// parameters.
class Variable : public Castable<Variable, Node> {
public:
/// Constructor for variables and non-overridable constants
/// Constructor
/// @param declaration the AST declaration node
/// @param type the variable type
/// @param storage_class the variable storage class
/// @param access the variable access control type
/// @param binding_point the optional resource binding point of the variable
Variable(const ast::Variable* declaration,
const sem::Type* type,
ast::StorageClass storage_class,
ast::Access access,
sem::BindingPoint binding_point = {});
/// Constructor for overridable pipeline constants
/// @param declaration the AST declaration node
/// @param type the variable type
/// @param constant_id the pipeline constant ID
Variable(const ast::Variable* declaration,
const sem::Type* type,
uint16_t constant_id);
ast::Access access);
/// Destructor
~Variable() override;
@ -74,32 +66,109 @@ class Variable : public Castable<Variable, Node> {
/// @returns the access control for the variable
ast::Access Access() const { return access_; }
/// @returns the resource binding point for the variable
sem::BindingPoint BindingPoint() const { return binding_point_; }
/// @returns the expressions that use the variable
const std::vector<const VariableUser*>& Users() const { return users_; }
/// @param user the user to add
void AddUser(const VariableUser* user) { users_.emplace_back(user); }
/// @returns true if this variable is an overridable pipeline constant
bool IsPipelineConstant() const { return is_pipeline_constant_; }
/// @returns the pipeline constant ID associated with the variable
uint16_t ConstantId() const { return constant_id_; }
private:
const ast::Variable* const declaration_;
const sem::Type* const type_;
ast::StorageClass const storage_class_;
ast::Access const access_;
sem::BindingPoint binding_point_;
std::vector<const VariableUser*> users_;
const bool is_pipeline_constant_;
const uint16_t constant_id_ = 0;
};
/// LocalVariable is a function-scope variable
class LocalVariable : public Castable<LocalVariable, Variable> {
public:
/// Constructor
/// @param declaration the AST declaration node
/// @param type the variable type
/// @param storage_class the variable storage class
/// @param access the variable access control type
LocalVariable(const ast::Variable* declaration,
const sem::Type* type,
ast::StorageClass storage_class,
ast::Access access);
/// Destructor
~LocalVariable() override;
};
/// GlobalVariable is a module-scope variable
class GlobalVariable : public Castable<GlobalVariable, Variable> {
public:
/// Constructor for non-overridable constants
/// @param declaration the AST declaration node
/// @param type the variable type
/// @param storage_class the variable storage class
/// @param access the variable access control type
/// @param binding_point the optional resource binding point of the variable
GlobalVariable(const ast::Variable* declaration,
const sem::Type* type,
ast::StorageClass storage_class,
ast::Access access,
sem::BindingPoint binding_point = {});
/// Constructor for overridable pipeline constants
/// @param declaration the AST declaration node
/// @param type the variable type
/// @param constant_id the pipeline constant ID
GlobalVariable(const ast::Variable* declaration,
const sem::Type* type,
uint16_t constant_id);
/// Destructor
~GlobalVariable() override;
/// @returns the resource binding point for the variable
sem::BindingPoint BindingPoint() const { return binding_point_; }
/// @returns the pipeline constant ID associated with the variable
uint16_t ConstantId() const { return constant_id_; }
/// @returns true if this variable is an overridable pipeline constant
bool IsPipelineConstant() const { return is_pipeline_constant_; }
private:
sem::BindingPoint binding_point_;
bool const is_pipeline_constant_;
uint16_t const constant_id_ = 0;
};
/// Parameter is a function parameter
class Parameter : public Castable<Parameter, Variable> {
public:
/// Constructor for function parameters
/// @param declaration the AST declaration node
/// @param type the variable type
/// @param storage_class the variable storage class
/// @param access the variable access control type
/// @param usage the semantic usage for the parameter
Parameter(const ast::Variable* declaration,
const sem::Type* type,
ast::StorageClass storage_class,
ast::Access access,
const ParameterUsage usage = ParameterUsage::kNone);
/// Copy constructor
Parameter(const Parameter&);
/// Destructor
~Parameter() override;
/// @returns the semantic usage for the parameter
ParameterUsage Usage() const { return usage_; }
private:
ParameterUsage const usage_;
};
/// ParameterList is a list of Parameter
using ParameterList = std::vector<const Parameter*>;
/// VariableUser holds the semantic information for an identifier expression
/// node that resolves to a variable.
class VariableUser : public Castable<VariableUser, Expression> {

8
src/transform/array_length_from_uniform.cc
View File

@ -131,7 +131,13 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx,
}
// Get the index to use for the buffer size array.
auto binding = storage_buffer_sem->Variable()->BindingPoint();
auto* var = tint::As<sem::GlobalVariable>(storage_buffer_sem->Variable());
if (!var) {
TINT_ICE(Transform, ctx.dst->Diagnostics())
<< "storage buffer is not a global variable";
break;
}
auto binding = var->BindingPoint();
auto idx_itr = cfg->bindpoint_to_size_index.find(binding);
if (idx_itr == cfg->bindpoint_to_size_index.end()) {
ctx.dst->Diagnostics().add_error(

4
src/transform/decompose_memory_access.cc
View File

@ -682,8 +682,8 @@ struct DecomposeMemoryAccess::State {
// Other parameters are copied as-is:
for (size_t i = 1; i < intrinsic->Parameters().size(); i++) {
auto& param = intrinsic->Parameters()[i];
auto* ty = CreateASTTypeFor(ctx, param.type);
auto* param = intrinsic->Parameters()[i];
auto* ty = CreateASTTypeFor(ctx, param->Type());
params.emplace_back(b.Param("param_" + std::to_string(i), ty));
}

6
src/transform/msl.cc
View File

@ -66,10 +66,10 @@ Output Msl::Run(const Program* in, const DataMap& inputs) {
// Use the SSBO binding numbers as the indices for the buffer size lookups.
for (auto* var : in->AST().GlobalVariables()) {
auto* sem_var = in->Sem().Get(var);
if (sem_var->StorageClass() == ast::StorageClass::kStorage) {
auto* global = in->Sem().Get<sem::GlobalVariable>(var);
if (global && global->StorageClass() == ast::StorageClass::kStorage) {
array_length_from_uniform_cfg.bindpoint_to_size_index.emplace(
sem_var->BindingPoint(), sem_var->BindingPoint().binding);
global->BindingPoint(), global->BindingPoint().binding);
}
}

2
src/transform/robustness.cc
View File

@ -219,7 +219,7 @@ struct Robustness::State {
auto* texture_arg = expr->params()[texture_idx];
auto* coords_arg = expr->params()[coords_idx];
auto* coords_ty = intrinsic->Parameters()[coords_idx].type;
auto* coords_ty = intrinsic->Parameters()[coords_idx]->Type();
// If the level is provided, then we need to clamp this. As the level is
// used by textureDimensions() and the texture[Load|Store]() calls, we need

1
src/transform/transform.cc
View File

@ -22,6 +22,7 @@
#include "src/sem/block_statement.h"
#include "src/sem/for_loop_statement.h"
#include "src/sem/reference_type.h"
#include "src/sem/sampler_type.h"
TINT_INSTANTIATE_TYPEINFO(tint::transform::Transform);
TINT_INSTANTIATE_TYPEINFO(tint::transform::Data);

15
src/utils/hash.h
View File

@ -43,16 +43,16 @@ struct HashCombineOffset<8> {
static constexpr inline uint64_t value() { return 0x9e3779b97f4a7c16; }
};
// When hashing sparse structures we want to iteratively build a hash value with
// only parts of the data. HashCombine "hashes" together an existing hash and
// hashable values.
} // namespace detail
/// HashCombine "hashes" together an existing hash and hashable values.
template <typename T>
void HashCombine(size_t* hash, const T& value) {
constexpr size_t offset = HashCombineOffset<sizeof(size_t)>::value();
constexpr size_t offset = detail::HashCombineOffset<sizeof(size_t)>::value();
*hash ^= std::hash<T>()(value) + offset + (*hash << 6) + (*hash >> 2);
}
// Helper for hashing vectors
/// HashCombine "hashes" together an existing hash and hashable values.
template <typename T>
void HashCombine(size_t* hash, const std::vector<T>& vector) {
HashCombine(hash, vector.size());
@ -61,20 +61,19 @@ void HashCombine(size_t* hash, const std::vector<T>& vector) {
}
}
/// HashCombine "hashes" together an existing hash and hashable values.
template <typename T, typename... ARGS>
void HashCombine(size_t* hash, const T& value, const ARGS&... args) {
HashCombine(hash, value);
HashCombine(hash, args...);
}
} // namespace detail
/// @returns a hash of the combined arguments. The returned hash is dependent on
/// the order of the arguments.
template <typename... ARGS>
size_t Hash(const ARGS&... args) {
size_t hash = 102931; // seed with an arbitrary prime
detail::HashCombine(&hash, args...);
HashCombine(&hash, args...);
return hash;
}

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

@ -1143,7 +1143,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
// InterlockedExchange and discard the returned value
{ // T result = 0;
auto pre = line();
auto* value_ty = intrinsic->Parameters()[1].type;
auto* value_ty = intrinsic->Parameters()[1]->Type();
if (!EmitTypeAndName(pre, value_ty, ast::StorageClass::kNone,
ast::Access::kUndefined, result)) {
return false;
@ -1278,9 +1278,9 @@ bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
return CallIntrinsicHelper(
out, expr, intrinsic,
[&](TextBuffer* b, const std::vector<std::string>& params) {
auto* significand_ty = intrinsic->Parameters()[0].type;
auto* significand_ty = intrinsic->Parameters()[0]->Type();
auto significand = params[0];
auto* exponent_ty = intrinsic->Parameters()[1].type;
auto* exponent_ty = intrinsic->Parameters()[1]->Type();
auto exponent = params[1];
std::string width;
@ -1314,7 +1314,7 @@ bool GeneratorImpl::EmitIsNormalCall(std::ostream& out,
return CallIntrinsicHelper(
out, expr, intrinsic,
[&](TextBuffer* b, const std::vector<std::string>& params) {
auto* input_ty = intrinsic->Parameters()[0].type;
auto* input_ty = intrinsic->Parameters()[0]->Type();
std::string width;
if (auto* vec = input_ty->As<sem::Vector>()) {
@ -2434,12 +2434,13 @@ bool GeneratorImpl::EmitEntryPointFunction(ast::Function* func) {
}
if (wgsize[i].overridable_const) {
auto* sem_const = builder_.Sem().Get(wgsize[i].overridable_const);
if (!sem_const->IsPipelineConstant()) {
auto* global = builder_.Sem().Get<sem::GlobalVariable>(
wgsize[i].overridable_const);
if (!global->IsPipelineConstant()) {
TINT_ICE(Writer, builder_.Diagnostics())
<< "expected a pipeline-overridable constant";
}
out << kSpecConstantPrefix << sem_const->ConstantId();
out << kSpecConstantPrefix << global->ConstantId();
} else {
out << std::to_string(wgsize[i].value);
}
@ -3187,8 +3188,9 @@ bool GeneratorImpl::EmitProgramConstVariable(const ast::Variable* var) {
auto* sem = builder_.Sem().Get(var);
auto* type = sem->Type();
if (sem->IsPipelineConstant()) {
auto const_id = sem->ConstantId();
auto* global = sem->As<sem::GlobalVariable>();
if (global && global->IsPipelineConstant()) {
auto const_id = global->ConstantId();
line() << "#ifndef " << kSpecConstantPrefix << const_id;
@ -3247,12 +3249,12 @@ bool GeneratorImpl::CallIntrinsicHelper(std::ostream& out,
}
{
ScopedParen sp(decl);
for (auto param : intrinsic->Parameters()) {
for (auto* param : intrinsic->Parameters()) {
if (!parameter_names.empty()) {
decl << ", ";
}
auto param_name = "param_" + std::to_string(parameter_names.size());
const auto* ty = param.type;
const auto* ty = param->Type();
if (auto* ptr = ty->As<sem::Pointer>()) {
decl << "inout ";
ty = ptr->StoreType();

1
src/writer/hlsl/generator_impl_type_test.cc
View File

@ -19,6 +19,7 @@
#include "src/sem/depth_texture_type.h"
#include "src/sem/multisampled_texture_type.h"
#include "src/sem/sampled_texture_type.h"
#include "src/sem/sampler_type.h"
#include "src/sem/storage_texture_type.h"
#include "src/writer/hlsl/test_helper.h"

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

@ -2323,9 +2323,9 @@ bool GeneratorImpl::EmitProgramConstVariable(const ast::Variable* var) {
out << " " << program_->Symbols().NameFor(var->symbol());
}
auto* sem_var = program_->Sem().Get(var);
if (sem_var->IsPipelineConstant()) {
out << " [[function_constant(" << sem_var->ConstantId() << ")]]";
auto* global = program_->Sem().Get<sem::GlobalVariable>(var);
if (global && global->IsPipelineConstant()) {
out << " [[function_constant(" << global->ConstantId() << ")]]";
} else if (var->constructor() != nullptr) {
out << " = ";
if (!EmitExpression(out, var->constructor())) {

1
src/writer/msl/generator_impl_type_test.cc
View File

@ -20,6 +20,7 @@
#include "src/sem/depth_texture_type.h"
#include "src/sem/multisampled_texture_type.h"
#include "src/sem/sampled_texture_type.h"
#include "src/sem/sampler_type.h"
#include "src/sem/storage_texture_type.h"
#include "src/writer/msl/test_helper.h"

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

@ -492,7 +492,8 @@ bool Builder::GenerateExecutionModes(ast::Function* func, uint32_t id) {
auto constant = ScalarConstant::U32(wgsize[i].value);
if (wgsize[i].overridable_const) {
// Make the constant specializable.
auto* sem_const = builder_.Sem().Get(wgsize[i].overridable_const);
auto* sem_const = builder_.Sem().Get<sem::GlobalVariable>(
wgsize[i].overridable_const);
if (!sem_const->IsPipelineConstant()) {
TINT_ICE(Writer, builder_.Diagnostics())
<< "expected a pipeline-overridable constant";
@ -1635,10 +1636,10 @@ uint32_t Builder::GenerateLiteralIfNeeded(ast::Variable* var,
ast::Literal* lit) {
ScalarConstant constant;
auto* sem_var = builder_.Sem().Get(var);
if (sem_var && sem_var->IsPipelineConstant()) {
auto* global = builder_.Sem().Get<sem::GlobalVariable>(var);
if (global && global->IsPipelineConstant()) {
constant.is_spec_op = true;
constant.constant_id = sem_var->ConstantId();
constant.constant_id = global->ConstantId();
}
if (auto* l = lit->As<ast::BoolLiteral>()) {
@ -2295,13 +2296,13 @@ uint32_t Builder::GenerateIntrinsic(ast::CallExpression* call,
// and loads it if necessary. Returns 0 on error.
auto get_param_as_value_id = [&](size_t i) -> uint32_t {
auto* arg = call->params()[i];
auto& param = intrinsic->Parameters()[i];
auto* param = intrinsic->Parameters()[i];
auto val_id = GenerateExpression(arg);
if (val_id == 0) {
return 0;
}
if (!param.type->Is<sem::Pointer>()) {
if (!param->Type()->Is<sem::Pointer>()) {
val_id = GenerateLoadIfNeeded(TypeOf(arg), val_id);
}
return val_id;
@ -2527,7 +2528,8 @@ uint32_t Builder::GenerateIntrinsic(ast::CallExpression* call,
// splat the condition into a vector of the same size.
// TODO(jrprice): If we're targeting SPIR-V 1.4, we don't need to do this.
auto* result_vector_type = intrinsic->ReturnType()->As<sem::Vector>();
if (result_vector_type && intrinsic->Parameters()[2].type->is_scalar()) {
if (result_vector_type &&
intrinsic->Parameters()[2]->Type()->is_scalar()) {
sem::Bool bool_type;
sem::Vector bool_vec_type(&bool_type, result_vector_type->size());
if (!GenerateTypeIfNeeded(&bool_vec_type)) {
@ -3036,14 +3038,15 @@ bool Builder::GenerateAtomicIntrinsic(ast::CallExpression* call,
Operand result_type,
Operand result_id) {
auto is_value_signed = [&] {
return intrinsic->Parameters()[1].type->Is<sem::I32>();
return intrinsic->Parameters()[1]->Type()->Is<sem::I32>();
};
auto storage_class =
intrinsic->Parameters()[0].type->As<sem::Pointer>()->StorageClass();
intrinsic->Parameters()[0]->Type()->As<sem::Pointer>()->StorageClass();
uint32_t memory_id = 0;
switch (intrinsic->Parameters()[0].type->As<sem::Pointer>()->StorageClass()) {
switch (
intrinsic->Parameters()[0]->Type()->As<sem::Pointer>()->StorageClass()) {
case ast::StorageClass::kWorkgroup:
memory_id = GenerateConstantIfNeeded(
ScalarConstant::U32(static_cast<uint32_t>(spv::Scope::Workgroup)));

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

@ -35,6 +35,7 @@
#include "src/ast/variable_decl_statement.h"
#include "src/program_builder.h"
#include "src/scope_stack.h"
#include "src/sem/intrinsic.h"
#include "src/sem/storage_texture_type.h"
#include "src/writer/spirv/function.h"
#include "src/writer/spirv/scalar_constant.h"