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>
diff --git a/src/ast/identifier_expression.h b/src/ast/identifier_expression.h
index da32342..3deff16 100644
--- a/src/ast/identifier_expression.h
+++ b/src/ast/identifier_expression.h
@@ -16,7 +16,6 @@
#define SRC_AST_IDENTIFIER_EXPRESSION_H_
#include "src/ast/expression.h"
-#include "src/sem/intrinsic.h"
namespace tint {
namespace ast {
diff --git a/src/inspector/inspector.cc b/src/inspector/inspector.cc
index 1003e0f..0e64db7 100644
--- a/src/inspector/inspector.cc
+++ b/src/inspector/inspector.cc
@@ -174,7 +174,8 @@
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::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 @@
// 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<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;
diff --git a/src/inspector/inspector_test.cc b/src/inspector/inspector_test.cc
index 7d7740a..40fc591 100644
--- a/src/inspector/inspector_test.cc
+++ b/src/inspector/inspector_test.cc
@@ -928,16 +928,19 @@
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) {
diff --git a/src/intrinsic_table.cc b/src/intrinsic_table.cc
index 233aff4..4f1fa28 100644
--- a/src/intrinsic_table.cc
+++ b/src/intrinsic_table.cc
@@ -701,6 +701,55 @@
#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 @@
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 @@
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 @@
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 @@
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);
});
}
diff --git a/src/intrinsic_table_test.cc b/src/intrinsic_table_test.cc
index 2635312..374c56b 100644
--- a/src/intrinsic_table_test.cc
+++ b/src/intrinsic_table_test.cc
@@ -45,7 +45,8 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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) {
diff --git a/src/program_builder.h b/src/program_builder.h
index d4a25cf..a7d8f90 100644
--- a/src/program_builder.h
+++ b/src/program_builder.h
@@ -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"
diff --git a/src/reader/spirv/function.cc b/src/reader/spirv/function.cc
index d673c60..819b9a3 100644
--- a/src/reader/spirv/function.cc
+++ b/src/reader/spirv/function.cc
@@ -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 @@
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));
diff --git a/src/resolver/intrinsic_test.cc b/src/resolver/intrinsic_test.cc
index 27b15df..d94ed2d 100644
--- a/src/resolver/intrinsic_test.cc
+++ b/src/resolver/intrinsic_test.cc
@@ -1706,11 +1706,11 @@
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 << ")";
diff --git a/src/resolver/pipeline_overridable_constant_test.cc b/src/resolver/pipeline_overridable_constant_test.cc
index 9fdfd8f..5d0571f 100644
--- a/src/resolver/pipeline_overridable_constant_test.cc
+++ b/src/resolver/pipeline_overridable_constant_test.cc
@@ -30,7 +30,7 @@
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 @@
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 @@
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 @@
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()));
}
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index b41b1bc..842b348 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -3693,11 +3693,24 @@
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 @@
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);
diff --git a/src/resolver/resolver_test.cc b/src/resolver/resolver_test.cc
index d231774..f176dc1 100644
--- a/src/resolver/resolver_test.cc
+++ b/src/resolver/resolver_test.cc
@@ -2018,8 +2018,10 @@
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) {
diff --git a/src/sem/call_target.cc b/src/sem/call_target.cc
index 5306bf8..b8bf12c 100644
--- a/src/sem/call_target.cc
+++ b/src/sem/call_target.cc
@@ -32,18 +32,12 @@
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
diff --git a/src/sem/call_target.h b/src/sem/call_target.h
index e994c80..33d5029 100644
--- a/src/sem/call_target.h
+++ b/src/sem/call_target.h
@@ -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 @@
} // 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_
diff --git a/src/sem/function.cc b/src/sem/function.cc
index fcf192a..21fabf9 100644
--- a/src/sem/function.cc
+++ b/src/sem/function.cc
@@ -27,31 +27,17 @@
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)),
diff --git a/src/sem/function.h b/src/sem/function.h
index d48047b..6ab1d39 100644
--- a/src/sem/function.h
+++ b/src/sem/function.h
@@ -68,7 +68,7 @@
/// @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 @@
/// @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 @@
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_;
diff --git a/src/sem/intrinsic.cc b/src/sem/intrinsic.cc
index 4be75e8..b7d57a9 100644
--- a/src/sem/intrinsic.cc
+++ b/src/sem/intrinsic.cc
@@ -19,11 +19,6 @@
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 @@
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 @@
supported_stages_(supported_stages),
is_deprecated_(is_deprecated) {}
-Intrinsic::Intrinsic(const Intrinsic&) = default;
-
Intrinsic::~Intrinsic() = default;
bool Intrinsic::IsCoarseDerivative() const {
@@ -155,25 +148,5 @@
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
diff --git a/src/sem/intrinsic.h b/src/sem/intrinsic.h
index 25c0c59..74f332f 100644
--- a/src/sem/intrinsic.h
+++ b/src/sem/intrinsic.h
@@ -88,13 +88,10 @@
/// 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 @@
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
diff --git a/src/sem/intrinsic_type.cc b/src/sem/intrinsic_type.cc
index 75616ca..a9cac51 100644
--- a/src/sem/intrinsic_type.cc
+++ b/src/sem/intrinsic_type.cc
@@ -24,6 +24,8 @@
#include "src/sem/intrinsic_type.h"
+#include <sstream>
+
namespace tint {
namespace sem {
@@ -529,5 +531,10 @@
return "<unknown>";
}
+std::ostream& operator<<(std::ostream& out, IntrinsicType i) {
+ out << str(i);
+ return out;
+}
+
} // namespace sem
} // namespace tint
diff --git a/src/sem/intrinsic_type.cc.tmpl b/src/sem/intrinsic_type.cc.tmpl
index e4059e3..f02c13f 100644
--- a/src/sem/intrinsic_type.cc.tmpl
+++ b/src/sem/intrinsic_type.cc.tmpl
@@ -10,6 +10,8 @@
#include "src/sem/intrinsic_type.h"
+#include <sstream>
+
namespace tint {
namespace sem {
@@ -34,5 +36,10 @@
return "<unknown>";
}
+std::ostream& operator<<(std::ostream& out, IntrinsicType i) {
+ out << str(i);
+ return out;
+}
+
} // namespace sem
} // namespace tint
diff --git a/src/sem/intrinsic_type.h b/src/sem/intrinsic_type.h
index 29e754b..6e47c3e 100644
--- a/src/sem/intrinsic_type.h
+++ b/src/sem/intrinsic_type.h
@@ -26,6 +26,7 @@
#define SRC_SEM_INTRINSIC_TYPE_H_
#include <string>
+#include <sstream>
namespace tint {
namespace sem {
@@ -143,6 +144,10 @@
/// 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
diff --git a/src/sem/intrinsic_type.h.tmpl b/src/sem/intrinsic_type.h.tmpl
index 8882645..deb935c 100644
--- a/src/sem/intrinsic_type.h.tmpl
+++ b/src/sem/intrinsic_type.h.tmpl
@@ -12,6 +12,7 @@
#define SRC_SEM_INTRINSIC_TYPE_H_
#include <string>
+#include <sstream>
namespace tint {
namespace sem {
@@ -34,6 +35,10 @@
/// 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
diff --git a/src/sem/sampler_type_test.cc b/src/sem/sampler_type_test.cc
index f27ac2e..9f2d5fc 100644
--- a/src/sem/sampler_type_test.cc
+++ b/src/sem/sampler_type_test.cc
@@ -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"
diff --git a/src/sem/variable.cc b/src/sem/variable.cc
index 490f890..fb6f48f 100644
--- a/src/sem/variable.cc
+++ b/src/sem/variable.cc
@@ -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 @@
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,
- const sem::Type* type,
- uint16_t constant_id)
- : declaration_(declaration),
- type_(type),
- storage_class_(ast::StorageClass::kNone),
- access_(ast::Access::kReadWrite),
+GlobalVariable::GlobalVariable(const ast::Variable* declaration,
+ const sem::Type* type,
+ uint16_t constant_id)
+ : 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,
diff --git a/src/sem/variable.h b/src/sem/variable.h
index 356ed9b..fc8d9be 100644
--- a/src/sem/variable.h
+++ b/src/sem/variable.h
@@ -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 @@
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 @@
/// @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> {
diff --git a/src/transform/array_length_from_uniform.cc b/src/transform/array_length_from_uniform.cc
index 35fbcb9..8a0fa17 100644
--- a/src/transform/array_length_from_uniform.cc
+++ b/src/transform/array_length_from_uniform.cc
@@ -131,7 +131,13 @@
}
// 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(
diff --git a/src/transform/decompose_memory_access.cc b/src/transform/decompose_memory_access.cc
index f06b3ae..0e2b245 100644
--- a/src/transform/decompose_memory_access.cc
+++ b/src/transform/decompose_memory_access.cc
@@ -682,8 +682,8 @@
// 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));
}
diff --git a/src/transform/msl.cc b/src/transform/msl.cc
index 53ead18..5c86a03 100644
--- a/src/transform/msl.cc
+++ b/src/transform/msl.cc
@@ -66,10 +66,10 @@
// 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);
}
}
diff --git a/src/transform/robustness.cc b/src/transform/robustness.cc
index 92abd78..1e8b94e 100644
--- a/src/transform/robustness.cc
+++ b/src/transform/robustness.cc
@@ -219,7 +219,7 @@
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
diff --git a/src/transform/transform.cc b/src/transform/transform.cc
index cd788c5..f67c046 100644
--- a/src/transform/transform.cc
+++ b/src/transform/transform.cc
@@ -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);
diff --git a/src/utils/hash.h b/src/utils/hash.h
index bd78942..91792bf 100644
--- a/src/utils/hash.h
+++ b/src/utils/hash.h
@@ -43,16 +43,16 @@
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 @@
}
}
+/// 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;
}
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index fc9e505..65b44c9 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -1143,7 +1143,7 @@
// 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 @@
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 @@
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 @@
}
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 @@
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 @@
}
{
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();
diff --git a/src/writer/hlsl/generator_impl_type_test.cc b/src/writer/hlsl/generator_impl_type_test.cc
index feedf4e..7db8426 100644
--- a/src/writer/hlsl/generator_impl_type_test.cc
+++ b/src/writer/hlsl/generator_impl_type_test.cc
@@ -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"
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index 39292bd..c22b146 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -2323,9 +2323,9 @@
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())) {
diff --git a/src/writer/msl/generator_impl_type_test.cc b/src/writer/msl/generator_impl_type_test.cc
index fa55b75..646ae7a 100644
--- a/src/writer/msl/generator_impl_type_test.cc
+++ b/src/writer/msl/generator_impl_type_test.cc
@@ -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"
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index d9bf363..f8e8ead 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -492,7 +492,8 @@
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 @@
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 @@
// 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 @@
// 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 @@
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)));
diff --git a/src/writer/spirv/builder.h b/src/writer/spirv/builder.h
index a81da1e..c33db7c 100644
--- a/src/writer/spirv/builder.h
+++ b/src/writer/spirv/builder.h
@@ -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"
