Import Tint changes from Dawn
Changes:
- d3f2b788cf964497ed3f01ab19359e76fcc2fc06 [tint][ir] Add Block() method to Instruction by Ben Clayton <bclayton@google.com>
- 293b4a5fc74eb7ba720e50b6138b2765538ee6cb [tint][ir] from_program.cc: Rename current_flow_block_ by Ben Clayton <bclayton@google.com>
- 212959b7383a69d4df61a17bc5710396411d060c [ir] Change function return attributes to not be list. by dan sinclair <dsinclair@chromium.org>
- 0df9b0312c1d2744125d3ff96bfdb856a3ca87b8 [ir] Add binding point to global vars. by dan sinclair <dsinclair@chromium.org>
- 765d45f31c086762ba480e4f33b6532c4cafc85f [ir] Add parameter attributes. by dan sinclair <dsinclair@chromium.org>
GitOrigin-RevId: d3f2b788cf964497ed3f01ab19359e76fcc2fc06
Change-Id: Iab919cf6bc4e3b0c35f563b7a6301109935cc875
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/134884
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index bc74d5d..83fcbf5 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1209,6 +1209,7 @@
sources = [
"ir/binary.cc",
"ir/binary.h",
+ "ir/binding_point.h",
"ir/bitcast.cc",
"ir/bitcast.h",
"ir/block.cc",
@@ -1253,6 +1254,7 @@
"ir/instruction.h",
"ir/load.cc",
"ir/load.h",
+ "ir/location.h",
"ir/loop.cc",
"ir/loop.h",
"ir/module.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 277d819..c20c850 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -716,6 +716,7 @@
list(APPEND TINT_LIB_SRCS
ir/binary.cc
ir/binary.h
+ ir/binding_point.h
ir/bitcast.cc
ir/bitcast.h
ir/block.cc
@@ -762,6 +763,7 @@
ir/instruction.h
ir/load.cc
ir/load.h
+ ir/location.h
ir/loop.cc
ir/loop.h
ir/module.cc
diff --git a/src/tint/ir/binding_point.h b/src/tint/ir/binding_point.h
new file mode 100644
index 0000000..9db225e
--- /dev/null
+++ b/src/tint/ir/binding_point.h
@@ -0,0 +1,32 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_IR_BINDING_POINT_H_
+#define SRC_TINT_IR_BINDING_POINT_H_
+
+#include <cstdint>
+
+namespace tint::ir {
+
+/// Binding information
+struct BindingPoint {
+ /// The `@group` part of the binding point
+ uint32_t group = 0;
+ /// The `@binding` part of the binding point
+ uint32_t binding = 0;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_BINDING_POINT_H_
diff --git a/src/tint/ir/block.cc b/src/tint/ir/block.cc
index 2030c19..4da1287 100644
--- a/src/tint/ir/block.cc
+++ b/src/tint/ir/block.cc
@@ -22,4 +22,16 @@
Block::~Block() = default;
+void Block::AddInstruction(Instruction* inst) {
+ instructions_.Push(inst);
+ inst->SetBlock(this);
+}
+
+void Block::SetInstructions(utils::VectorRef<Instruction*> instructions) {
+ for (auto* i : instructions) {
+ i->SetBlock(this);
+ }
+ instructions_ = std::move(instructions);
+}
+
} // namespace tint::ir
diff --git a/src/tint/ir/block.h b/src/tint/ir/block.h
index d40002d..20a2da7 100644
--- a/src/tint/ir/block.h
+++ b/src/tint/ir/block.h
@@ -48,14 +48,14 @@
/// Sets the instructions in the block
/// @param instructions the instructions to set
- void SetInstructions(utils::VectorRef<const Instruction*> instructions) {
- instructions_ = std::move(instructions);
- }
+ void SetInstructions(utils::VectorRef<Instruction*> instructions);
/// @returns the instructions in the block
utils::VectorRef<const Instruction*> Instructions() const { return instructions_; }
- /// @returns the instructions in the block
- utils::Vector<const Instruction*, 16>& Instructions() { return instructions_; }
+
+ /// Adds the instruction to the end of the block
+ /// @param inst the instruction to add
+ void AddInstruction(Instruction* inst);
/// Sets the params to the block
/// @param params the params for the block
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 75cc2aa..d8e472a 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -139,39 +139,119 @@
Indent() << "}" << std::endl;
}
+void Disassembler::EmitBindingPoint(BindingPoint p) {
+ out_ << "@binding_point(" << p.group << ", " << p.binding << ")";
+}
+
+void Disassembler::EmitLocation(Location loc) {
+ out_ << "@location(" << loc.value << ")";
+ if (loc.interpolation.has_value()) {
+ out_ << ", @interpolate(";
+ out_ << loc.interpolation->type;
+ if (loc.interpolation->sampling != builtin::InterpolationSampling::kUndefined) {
+ out_ << ", ";
+ out_ << loc.interpolation->sampling;
+ }
+ out_ << ")";
+ }
+}
+
+void Disassembler::EmitParamAttributes(const FunctionParam* p) {
+ if (!p->Invariant() && !p->Location().has_value() && !p->BindingPoint().has_value() &&
+ !p->Builtin().has_value()) {
+ return;
+ }
+
+ out_ << " [";
+
+ bool need_comma = false;
+ auto comma = [&]() {
+ if (need_comma) {
+ out_ << ", ";
+ }
+ };
+
+ if (p->Invariant()) {
+ comma();
+ out_ << "@invariant";
+ need_comma = true;
+ }
+ if (p->Location().has_value()) {
+ EmitLocation(p->Location().value());
+ need_comma = true;
+ }
+ if (p->BindingPoint().has_value()) {
+ comma();
+ EmitBindingPoint(p->BindingPoint().value());
+ need_comma = true;
+ }
+ if (p->Builtin().has_value()) {
+ comma();
+ out_ << "@" << p->Builtin().value();
+ need_comma = true;
+ }
+ out_ << "]";
+}
+
+void Disassembler::EmitReturnAttributes(const Function* func) {
+ if (!func->ReturnInvariant() && !func->ReturnLocation().has_value() &&
+ !func->ReturnBuiltin().has_value()) {
+ return;
+ }
+
+ out_ << " [";
+
+ bool need_comma = false;
+ auto comma = [&]() {
+ if (need_comma) {
+ out_ << ", ";
+ }
+ };
+ if (func->ReturnInvariant()) {
+ comma();
+ out_ << "@invariant";
+ need_comma = true;
+ }
+ if (func->ReturnLocation().has_value()) {
+ comma();
+ EmitLocation(func->ReturnLocation().value());
+ need_comma = true;
+ }
+ if (func->ReturnBuiltin().has_value()) {
+ comma();
+ out_ << "@" << func->ReturnBuiltin().value();
+ need_comma = true;
+ }
+ out_ << "]";
+}
+
void Disassembler::EmitFunction(const Function* func) {
in_function_ = true;
- Indent() << "%" << IdOf(func) << " = func(";
- for (auto* p : func->Params()) {
+ Indent() << "%" << IdOf(func) << " =";
+
+ if (func->Stage() != Function::PipelineStage::kUndefined) {
+ out_ << " @" << func->Stage();
+ }
+ if (func->WorkgroupSize()) {
+ auto arr = func->WorkgroupSize().value();
+ out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2] << ")";
+ }
+
+ out_ << " func(";
+
+ for (const auto* p : func->Params()) {
if (p != func->Params().Front()) {
out_ << ", ";
}
out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
+
+ EmitParamAttributes(p);
}
out_ << "):" << func->ReturnType()->FriendlyName();
- if (func->Stage() != Function::PipelineStage::kUndefined) {
- out_ << " [@" << func->Stage();
+ EmitReturnAttributes(func);
- if (func->WorkgroupSize()) {
- auto arr = func->WorkgroupSize().value();
- out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2] << ")";
- }
-
- if (!func->ReturnAttributes().IsEmpty()) {
- out_ << " ra:";
-
- for (auto attr : func->ReturnAttributes()) {
- out_ << " @" << attr;
- if (attr == Function::ReturnAttribute::kLocation) {
- out_ << "(" << func->ReturnLocation().value() << ")";
- }
- }
- }
-
- out_ << "]";
- }
out_ << " -> %b" << IdOf(func->StartTarget()) << " {" << std::endl;
{
@@ -302,6 +382,11 @@
out_ << ", ";
EmitValue(v->Initializer());
}
+ if (v->BindingPoint().has_value()) {
+ out_ << " ";
+ EmitBindingPoint(v->BindingPoint().value());
+ }
+
out_ << std::endl;
},
[&](const ir::Branch* b) { EmitBranch(b); },
diff --git a/src/tint/ir/disassembler.h b/src/tint/ir/disassembler.h
index 06f5b08..9194942 100644
--- a/src/tint/ir/disassembler.h
+++ b/src/tint/ir/disassembler.h
@@ -59,6 +59,10 @@
void Walk(const Block* blk);
void WalkInternal(const Block* blk);
void EmitFunction(const Function* func);
+ void EmitParamAttributes(const FunctionParam* p);
+ void EmitReturnAttributes(const Function* func);
+ void EmitBindingPoint(BindingPoint p);
+ void EmitLocation(Location loc);
void EmitInstruction(const Instruction* inst);
void EmitValueWithType(const Value* val);
void EmitValue(const Value* val);
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index 20555d0..3dc52e1 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -43,6 +43,7 @@
#include "src/tint/ast/if_statement.h"
#include "src/tint/ast/increment_decrement_statement.h"
#include "src/tint/ast/int_literal_expression.h"
+#include "src/tint/ast/interpolate_attribute.h"
#include "src/tint/ast/invariant_attribute.h"
#include "src/tint/ast/let.h"
#include "src/tint/ast/literal_expression.h"
@@ -140,8 +141,8 @@
/// The stack of control blocks.
utils::Vector<Branch*, 8> control_stack_;
- /// The current flow block for expressions.
- Block* current_flow_block_ = nullptr;
+ /// The current block for expressions.
+ Block* current_block_ = nullptr;
/// The current function being processed.
Function* current_function_ = nullptr;
@@ -166,14 +167,14 @@
diagnostics_.add_error(tint::diag::System::IR, err, s);
}
- bool NeedBranch() { return current_flow_block_ && !current_flow_block_->HasBranchTarget(); }
+ bool NeedBranch() { return current_block_ && !current_block_->HasBranchTarget(); }
void SetBranch(Branch* br) {
- TINT_ASSERT(IR, current_flow_block_);
- TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget());
+ TINT_ASSERT(IR, current_block_);
+ TINT_ASSERT(IR, !current_block_->HasBranchTarget());
- current_flow_block_->Instructions().Push(br);
- current_flow_block_ = nullptr;
+ current_block_->AddInstruction(br);
+ current_block_ = nullptr;
}
Branch* FindEnclosingControl(ControlFlags flags) {
@@ -207,7 +208,7 @@
[&](const ast::Variable* var) {
// Setup the current flow node to be the root block for the module. The builder
// will handle creating it if it doesn't exist already.
- TINT_SCOPED_ASSIGNMENT(current_flow_block_, builder_.CreateRootBlockIfNeeded());
+ TINT_SCOPED_ASSIGNMENT(current_block_, builder_.CreateRootBlockIfNeeded());
EmitVariable(var);
},
[&](const ast::Function* func) { EmitFunction(func); },
@@ -230,6 +231,24 @@
return ResultType{std::move(mod)};
}
+ builtin::Interpolation ExtractInterpolation(const ast::InterpolateAttribute* interp) {
+ auto type = program_->Sem()
+ .Get(interp->type)
+ ->As<sem::BuiltinEnumExpression<builtin::InterpolationType>>();
+ builtin::InterpolationType interpolation_type = type->Value();
+
+ builtin::InterpolationSampling interpolation_sampling =
+ builtin::InterpolationSampling::kUndefined;
+ if (interp->sampling) {
+ auto sampling = program_->Sem()
+ .Get(interp->sampling)
+ ->As<sem::BuiltinEnumExpression<builtin::InterpolationSampling>>();
+ interpolation_sampling = sampling->Value();
+ }
+
+ return builtin::Interpolation{interpolation_type, interpolation_sampling};
+ }
+
void EmitFunction(const ast::Function* ast_func) {
// The flow stack should have been emptied when the previous function finished building.
TINT_ASSERT(IR, control_stack_.IsEmpty());
@@ -265,16 +284,16 @@
}
}
- utils::Vector<Function::ReturnAttribute, 1> return_attributes;
+ // Note, interpolated is only valid when paired with Location, so it will only be set
+ // when the location is set.
+ std::optional<builtin::Interpolation> interpolation;
for (auto* attr : ast_func->return_type_attributes) {
tint::Switch(
attr, //
- [&](const ast::LocationAttribute*) {
- return_attributes.Push(Function::ReturnAttribute::kLocation);
+ [&](const ast::InterpolateAttribute* interp) {
+ interpolation = ExtractInterpolation(interp);
},
- [&](const ast::InvariantAttribute*) {
- return_attributes.Push(Function::ReturnAttribute::kInvariant);
- },
+ [&](const ast::InvariantAttribute*) { ir_func->SetReturnInvariant(true); },
[&](const ast::BuiltinAttribute* b) {
if (auto* ident_sem =
program_->Sem()
@@ -282,13 +301,13 @@
->As<sem::BuiltinEnumExpression<builtin::BuiltinValue>>()) {
switch (ident_sem->Value()) {
case builtin::BuiltinValue::kPosition:
- return_attributes.Push(Function::ReturnAttribute::kPosition);
+ ir_func->SetReturnBuiltin(Function::ReturnBuiltin::kPosition);
break;
case builtin::BuiltinValue::kFragDepth:
- return_attributes.Push(Function::ReturnAttribute::kFragDepth);
+ ir_func->SetReturnBuiltin(Function::ReturnBuiltin::kFragDepth);
break;
case builtin::BuiltinValue::kSampleMask:
- return_attributes.Push(Function::ReturnAttribute::kSampleMask);
+ ir_func->SetReturnBuiltin(Function::ReturnBuiltin::kSampleMask);
break;
default:
TINT_ICE(IR, diagnostics_)
@@ -302,26 +321,98 @@
}
});
}
- ir_func->SetReturnAttributes(return_attributes);
+ if (sem->ReturnLocation().has_value()) {
+ ir_func->SetReturnLocation(sem->ReturnLocation().value(), interpolation);
+ }
}
- ir_func->SetReturnLocation(sem->ReturnLocation());
scopes_.Push();
TINT_DEFER(scopes_.Pop());
utils::Vector<FunctionParam*, 1> params;
for (auto* p : ast_func->params) {
- const auto* param_sem = program_->Sem().Get(p);
+ const auto* param_sem = program_->Sem().Get(p)->As<sem::Parameter>();
auto* ty = param_sem->Type()->Clone(clone_ctx_.type_ctx);
auto* param = builder_.FunctionParam(ty);
+ // Note, interpolated is only valid when paired with Location, so it will only be set
+ // when the location is set.
+ std::optional<builtin::Interpolation> interpolation;
+ for (auto* attr : p->attributes) {
+ tint::Switch(
+ attr, //
+ [&](const ast::InterpolateAttribute* interp) {
+ interpolation = ExtractInterpolation(interp);
+ },
+ [&](const ast::InvariantAttribute*) { param->SetInvariant(true); },
+ [&](const ast::BuiltinAttribute* b) {
+ if (auto* ident_sem =
+ program_->Sem()
+ .Get(b)
+ ->As<sem::BuiltinEnumExpression<builtin::BuiltinValue>>()) {
+ switch (ident_sem->Value()) {
+ case builtin::BuiltinValue::kVertexIndex:
+ param->SetBuiltin(FunctionParam::Builtin::kVertexIndex);
+ break;
+ case builtin::BuiltinValue::kInstanceIndex:
+ param->SetBuiltin(FunctionParam::Builtin::kInstanceIndex);
+ break;
+ case builtin::BuiltinValue::kPosition:
+ param->SetBuiltin(FunctionParam::Builtin::kPosition);
+ break;
+ case builtin::BuiltinValue::kFrontFacing:
+ param->SetBuiltin(FunctionParam::Builtin::kFrontFacing);
+ break;
+ case builtin::BuiltinValue::kLocalInvocationId:
+ param->SetBuiltin(FunctionParam::Builtin::kLocalInvocationId);
+ break;
+ case builtin::BuiltinValue::kLocalInvocationIndex:
+ param->SetBuiltin(
+ FunctionParam::Builtin::kLocalInvocationIndex);
+ break;
+ case builtin::BuiltinValue::kGlobalInvocationId:
+ param->SetBuiltin(FunctionParam::Builtin::kGlobalInvocationId);
+ break;
+ case builtin::BuiltinValue::kWorkgroupId:
+ param->SetBuiltin(FunctionParam::Builtin::kWorkgroupId);
+ break;
+ case builtin::BuiltinValue::kNumWorkgroups:
+ param->SetBuiltin(FunctionParam::Builtin::kNumWorkgroups);
+ break;
+ case builtin::BuiltinValue::kSampleIndex:
+ param->SetBuiltin(FunctionParam::Builtin::kSampleIndex);
+ break;
+ case builtin::BuiltinValue::kSampleMask:
+ param->SetBuiltin(FunctionParam::Builtin::kSampleMask);
+ break;
+ default:
+ TINT_ICE(IR, diagnostics_)
+ << "Unknown builtin value in parameter attributes "
+ << ident_sem->Value();
+ return;
+ }
+ } else {
+ TINT_ICE(IR, diagnostics_) << "Builtin attribute sem invalid";
+ return;
+ }
+ });
+
+ if (param_sem->Location().has_value()) {
+ param->SetLocation(param_sem->Location().value(), interpolation);
+ }
+ if (param_sem->BindingPoint().has_value()) {
+ param->SetBindingPoint(param_sem->BindingPoint()->group,
+ param_sem->BindingPoint()->binding);
+ }
+ }
+
scopes_.Set(p->name->symbol, param);
builder_.ir.SetName(param, p->name->symbol.NameView());
params.Push(param);
}
ir_func->SetParams(params);
- current_flow_block_ = ir_func->StartTarget();
+ current_block_ = ir_func->StartTarget();
EmitBlock(ast_func->body);
// If the branch target has already been set then a `return` was called. Only set in
@@ -331,7 +422,7 @@
}
TINT_ASSERT(IR, control_stack_.IsEmpty());
- current_flow_block_ = nullptr;
+ current_block_ = nullptr;
current_function_ = nullptr;
}
@@ -396,7 +487,7 @@
return;
}
auto store = builder_.Store(lhs.Get(), rhs.Get());
- current_flow_block_->Instructions().Push(store);
+ current_block_->AddInstruction(store);
}
void EmitIncrementDecrement(const ast::IncrementDecrementStatement* stmt) {
@@ -407,7 +498,7 @@
// Load from the LHS.
auto* lhs_value = builder_.Load(lhs.Get());
- current_flow_block_->Instructions().Push(lhs_value);
+ current_block_->AddInstruction(lhs_value);
auto* ty = lhs_value->Type();
@@ -420,10 +511,10 @@
} else {
inst = builder_.Subtract(ty, lhs_value, rhs);
}
- current_flow_block_->Instructions().Push(inst);
+ current_block_->AddInstruction(inst);
auto store = builder_.Store(lhs.Get(), inst);
- current_flow_block_->Instructions().Push(store);
+ current_block_->AddInstruction(store);
}
void EmitCompoundAssignment(const ast::CompoundAssignmentStatement* stmt) {
@@ -439,7 +530,7 @@
// Load from the LHS.
auto* lhs_value = builder_.Load(lhs.Get());
- current_flow_block_->Instructions().Push(lhs_value);
+ current_block_->AddInstruction(lhs_value);
auto* ty = lhs_value->Type();
@@ -489,10 +580,10 @@
TINT_ICE(IR, diagnostics_) << "missing binary operand type";
return;
}
- current_flow_block_->Instructions().Push(inst);
+ current_block_->AddInstruction(inst);
auto store = builder_.Store(lhs.Get(), inst);
- current_flow_block_->Instructions().Push(store);
+ current_block_->AddInstruction(store);
}
void EmitBlock(const ast::BlockStatement* block) {
@@ -512,12 +603,12 @@
return;
}
auto* if_inst = builder_.CreateIf(reg.Get());
- current_flow_block_->Instructions().Push(if_inst);
+ current_block_->AddInstruction(if_inst);
{
ControlStackScope scope(this, if_inst);
- current_flow_block_ = if_inst->True();
+ current_block_ = if_inst->True();
EmitBlock(stmt->body);
// If the true branch did not execute control flow, then go to the Merge().target
@@ -525,7 +616,7 @@
SetBranch(builder_.ExitIf(if_inst));
}
- current_flow_block_ = if_inst->False();
+ current_block_ = if_inst->False();
if (stmt->else_statement) {
EmitStatement(stmt->else_statement);
}
@@ -535,23 +626,23 @@
SetBranch(builder_.ExitIf(if_inst));
}
}
- current_flow_block_ = nullptr;
+ current_block_ = nullptr;
// If both branches went somewhere, then they both returned, continued or broke. So,
// there is no need for the if merge-block and there is nothing to branch to the merge
// block anyway.
if (IsConnected(if_inst->Merge())) {
- current_flow_block_ = if_inst->Merge();
+ current_block_ = if_inst->Merge();
}
}
void EmitLoop(const ast::LoopStatement* stmt) {
auto* loop_inst = builder_.CreateLoop();
- current_flow_block_->Instructions().Push(loop_inst);
+ current_block_->AddInstruction(loop_inst);
{
ControlStackScope scope(this, loop_inst);
- current_flow_block_ = loop_inst->Start();
+ current_block_ = loop_inst->Start();
// The loop doesn't use EmitBlock because it needs the scope stack to not get popped
// until after the continuing block.
@@ -569,7 +660,7 @@
// continue so we have to set the current block and then emit the branch if needed
// below otherwise empty continuing blocks will fail to branch back to the start
// block.
- current_flow_block_ = loop_inst->Continuing();
+ current_block_ = loop_inst->Continuing();
if (stmt->continuing) {
EmitBlock(stmt->continuing);
}
@@ -583,24 +674,24 @@
// The loop merge can get disconnected if the loop returns directly, or the continuing
// target branches, eventually, to the merge, but nothing branched to the
// Continuing() block.
- current_flow_block_ = loop_inst->Merge();
+ current_block_ = loop_inst->Merge();
if (!IsConnected(loop_inst->Merge())) {
- current_flow_block_ = nullptr;
+ current_block_ = nullptr;
}
}
void EmitWhile(const ast::WhileStatement* stmt) {
auto* loop_inst = builder_.CreateLoop();
- current_flow_block_->Instructions().Push(loop_inst);
+ current_block_->AddInstruction(loop_inst);
// Continue is always empty, just go back to the start
- current_flow_block_ = loop_inst->Continuing();
+ current_block_ = loop_inst->Continuing();
SetBranch(builder_.NextIteration(loop_inst));
{
ControlStackScope scope(this, loop_inst);
- current_flow_block_ = loop_inst->Start();
+ current_block_ = loop_inst->Start();
// Emit the while condition into the Start().target of the loop
auto reg = EmitExpression(stmt->condition);
@@ -610,15 +701,15 @@
// Create an `if (cond) {} else {break;}` control flow
auto* if_inst = builder_.CreateIf(reg.Get());
- current_flow_block_->Instructions().Push(if_inst);
+ current_block_->AddInstruction(if_inst);
- current_flow_block_ = if_inst->True();
+ current_block_ = if_inst->True();
SetBranch(builder_.ExitIf(if_inst));
- current_flow_block_ = if_inst->False();
+ current_block_ = if_inst->False();
SetBranch(builder_.ExitLoop(loop_inst));
- current_flow_block_ = if_inst->Merge();
+ current_block_ = if_inst->Merge();
EmitBlock(stmt->body);
if (NeedBranch()) {
@@ -627,12 +718,12 @@
}
// The while loop always has a path to the Merge().target as the break statement comes
// before anything inside the loop.
- current_flow_block_ = loop_inst->Merge();
+ current_block_ = loop_inst->Merge();
}
void EmitForLoop(const ast::ForLoopStatement* stmt) {
auto* loop_inst = builder_.CreateLoop();
- current_flow_block_->Instructions().Push(loop_inst);
+ current_block_->AddInstruction(loop_inst);
// Make sure the initializer ends up in a contained scope
scopes_.Push();
@@ -646,7 +737,7 @@
{
ControlStackScope scope(this, loop_inst);
- current_flow_block_ = loop_inst->Start();
+ current_block_ = loop_inst->Start();
if (stmt->condition) {
// Emit the condition into the target target of the loop
@@ -657,15 +748,15 @@
// Create an `if (cond) {} else {break;}` control flow
auto* if_inst = builder_.CreateIf(reg.Get());
- current_flow_block_->Instructions().Push(if_inst);
+ current_block_->AddInstruction(if_inst);
- current_flow_block_ = if_inst->True();
+ current_block_ = if_inst->True();
SetBranch(builder_.ExitIf(if_inst));
- current_flow_block_ = if_inst->False();
+ current_block_ = if_inst->False();
SetBranch(builder_.ExitLoop(loop_inst));
- current_flow_block_ = if_inst->Merge();
+ current_block_ = if_inst->Merge();
}
EmitBlock(stmt->body);
@@ -674,7 +765,7 @@
}
if (stmt->continuing) {
- current_flow_block_ = loop_inst->Continuing();
+ current_block_ = loop_inst->Continuing();
EmitStatement(stmt->continuing);
SetBranch(builder_.NextIteration(loop_inst));
}
@@ -682,7 +773,7 @@
// The while loop always has a path to the Merge().target as the break statement comes
// before anything inside the loop.
- current_flow_block_ = loop_inst->Merge();
+ current_block_ = loop_inst->Merge();
}
void EmitSwitch(const ast::SwitchStatement* stmt) {
@@ -692,7 +783,7 @@
return;
}
auto* switch_inst = builder_.CreateSwitch(reg.Get());
- current_flow_block_->Instructions().Push(switch_inst);
+ current_block_->AddInstruction(switch_inst);
{
ControlStackScope scope(this, switch_inst);
@@ -708,7 +799,7 @@
}
}
- current_flow_block_ = builder_.CreateCase(switch_inst, selectors);
+ current_block_ = builder_.CreateCase(switch_inst, selectors);
EmitBlock(c->Body()->Declaration());
if (NeedBranch()) {
@@ -716,10 +807,10 @@
}
}
}
- current_flow_block_ = nullptr;
+ current_block_ = nullptr;
if (IsConnected(switch_inst->Merge())) {
- current_flow_block_ = switch_inst->Merge();
+ current_block_ = switch_inst->Merge();
}
}
@@ -765,7 +856,7 @@
// figuring out the multi-level exit that is triggered.
void EmitDiscard(const ast::DiscardStatement*) {
auto* inst = builder_.Discard();
- current_flow_block_->Instructions().Push(inst);
+ current_block_->AddInstruction(inst);
}
void EmitBreakIf(const ast::BreakIfStatement* stmt) {
@@ -823,7 +914,7 @@
// If this expression maps to sem::Load, insert a load instruction to get the result.
if (result && sem->Is<sem::Load>()) {
auto* load = builder_.Load(result.Get());
- current_flow_block_->Instructions().Push(load);
+ current_block_->AddInstruction(load);
return load;
}
@@ -849,7 +940,12 @@
}
val->SetInitializer(init.Get());
}
- current_flow_block_->Instructions().Push(val);
+ current_block_->AddInstruction(val);
+
+ if (auto* gv = sem->As<sem::GlobalVariable>(); gv && var->HasBindingPoint()) {
+ val->SetBindingPoint(gv->BindingPoint().value().group,
+ gv->BindingPoint().value().binding);
+ }
// Store the declaration so we can get the instruction to store too
scopes_.Set(v->name->symbol, val);
@@ -916,7 +1012,7 @@
break;
}
- current_flow_block_->Instructions().Push(inst);
+ current_block_->AddInstruction(inst);
return inst;
}
@@ -940,7 +1036,7 @@
}
auto* if_inst = builder_.CreateIf(lhs.Get());
- current_flow_block_->Instructions().Push(if_inst);
+ current_block_->AddInstruction(if_inst);
auto* result = builder_.BlockParam(builder_.ir.Types().bool_());
if_inst->Merge()->SetParams(utils::Vector{result});
@@ -957,17 +1053,17 @@
if (expr->op == ast::BinaryOp::kLogicalAnd) {
// If the lhs is false, then that is the result we want to pass to the merge
// block as our argument
- current_flow_block_ = if_inst->False();
+ current_block_ = if_inst->False();
SetBranch(builder_.ExitIf(if_inst, std::move(alt_args)));
- current_flow_block_ = if_inst->True();
+ current_block_ = if_inst->True();
} else {
// If the lhs is true, then that is the result we want to pass to the merge
// block as our argument
- current_flow_block_ = if_inst->True();
+ current_block_ = if_inst->True();
SetBranch(builder_.ExitIf(if_inst, std::move(alt_args)));
- current_flow_block_ = if_inst->False();
+ current_block_ = if_inst->False();
}
rhs = EmitExpression(expr->rhs);
@@ -979,7 +1075,7 @@
SetBranch(builder_.ExitIf(if_inst, std::move(args)));
}
- current_flow_block_ = if_inst->Merge();
+ current_block_ = if_inst->Merge();
return result;
}
@@ -1061,7 +1157,7 @@
return utils::Failure;
}
- current_flow_block_->Instructions().Push(inst);
+ current_block_->AddInstruction(inst);
return inst;
}
@@ -1075,7 +1171,7 @@
auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx);
auto* inst = builder_.Bitcast(ty, val.Get());
- current_flow_block_->Instructions().Push(inst);
+ current_block_->AddInstruction(inst);
return inst;
}
@@ -1139,7 +1235,7 @@
if (inst == nullptr) {
return utils::Failure;
}
- current_flow_block_->Instructions().Push(inst);
+ current_block_->AddInstruction(inst);
return inst;
}
diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc
index 8cc5604..5f7bcb2 100644
--- a/src/tint/ir/from_program_binary_test.cc
+++ b/src/tint/ir/from_program_binary_test.cc
@@ -39,7 +39,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = add %3, 4u
@@ -62,7 +62,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = add %3, 1u
@@ -86,7 +86,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = add %3, 1u
@@ -110,7 +110,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = sub %3, 4u
@@ -133,7 +133,7 @@
%v1:ptr<private, i32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:i32 = load %v1
%4:i32 = sub %3, 1i
@@ -157,7 +157,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = sub %3, 1u
@@ -181,7 +181,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = mul %3, 4u
@@ -204,7 +204,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = mul %3, 1u
@@ -228,7 +228,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = div %3, 4u
@@ -251,7 +251,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = div %3, 1u
@@ -275,7 +275,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = mod %3, 4u
@@ -298,7 +298,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = mod %3, 1u
@@ -322,7 +322,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = and %3, 4u
@@ -345,7 +345,7 @@
%v1:ptr<private, bool, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = load %v1
%4:bool = and %3, false
@@ -369,7 +369,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = or %3, 4u
@@ -392,7 +392,7 @@
%v1:ptr<private, bool, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = load %v1
%4:bool = or %3, false
@@ -416,7 +416,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = xor %3, 4u
@@ -439,7 +439,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = xor %3, 1u
@@ -463,7 +463,7 @@
ret true
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = call %my_func
if %3 [t: %b3, f: %b4, m: %b5]
@@ -515,7 +515,7 @@
ret true
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = call %my_func
if %3 [t: %b3, f: %b4, m: %b5]
@@ -567,7 +567,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = eq %3, 4u
@@ -590,7 +590,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = neq %3, 4u
@@ -613,7 +613,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = lt %3, 4u
@@ -636,7 +636,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = gt %3, 4u
@@ -659,7 +659,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = lte %3, 4u
@@ -682,7 +682,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:bool = gte %3, 4u
@@ -705,7 +705,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = shiftl %3, 4u
@@ -728,7 +728,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = shiftl %3, 1u
@@ -752,7 +752,7 @@
ret 0u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = shiftr %3, 4u
@@ -775,7 +775,7 @@
%v1:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = load %v1
%4:u32 = shiftr %3, 1u
@@ -801,7 +801,7 @@
ret 0.0f
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:f32 = call %my_func
%4:bool = lt %3, 2.0f
@@ -845,7 +845,7 @@
ret true
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%tint_symbol:bool = call %my_func, false
ret
diff --git a/src/tint/ir/from_program_builtin_test.cc b/src/tint/ir/from_program_builtin_test.cc
index acdbc8a..4d21b48 100644
--- a/src/tint/ir/from_program_builtin_test.cc
+++ b/src/tint/ir/from_program_builtin_test.cc
@@ -39,7 +39,7 @@
%i:ptr<private, f32, read_write> = var, 1.0f
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:f32 = load %i
%tint_symbol:f32 = asin %3
diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc
index 155e42c..d305855 100644
--- a/src/tint/ir/from_program_call_test.cc
+++ b/src/tint/ir/from_program_call_test.cc
@@ -40,7 +40,7 @@
ret 0.0f
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:f32 = call %my_func
%tint_symbol:f32 = bitcast %3
@@ -60,7 +60,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = func():void [@fragment] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = @fragment func():void -> %b1 {
%b1 = block {
discard
ret
@@ -82,7 +82,7 @@
ret
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%4:void = call %my_func, 6.0f
ret
@@ -104,7 +104,7 @@
%i:ptr<private, i32, read_write> = var, 1i
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:i32 = load %i
%tint_symbol:f32 = convert i32, %3
@@ -142,7 +142,7 @@
%i:ptr<private, f32, read_write> = var, 1.0f
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:f32 = load %i
%tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %3
diff --git a/src/tint/ir/from_program_function_test.cc b/src/tint/ir/from_program_function_test.cc
index 1be388a..57e78a3 100644
--- a/src/tint/ir/from_program_function_test.cc
+++ b/src/tint/ir/from_program_function_test.cc
@@ -34,7 +34,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec4<f32> [@vertex ra: @position] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test = @vertex func():vec4<f32> [@position] -> %b1 {
%b1 = block {
ret vec4<f32> 0.0f
}
@@ -49,7 +49,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():void [@fragment] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test = @fragment func():void -> %b1 {
%b1 = block {
ret
}
@@ -65,7 +65,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test = func():void [@compute @workgroup_size(8, 4, 2)] -> %b1 {
+ R"(%test = @compute @workgroup_size(8, 4, 2) func():void -> %b1 {
%b1 = block {
ret
}
@@ -96,7 +96,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec4<f32> [@vertex ra: @position] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test = @vertex func():vec4<f32> [@position] -> %b1 {
%b1 = block {
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
}
@@ -113,7 +113,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test = func():vec4<f32> [@vertex ra: @position @invariant] -> %b1 {
+ R"(%test = @vertex func():vec4<f32> [@invariant, @position] -> %b1 {
%b1 = block {
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
}
@@ -129,7 +129,26 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test = func():vec4<f32> [@fragment ra: @location(1)] -> %b1 {
+ R"(%test = @fragment func():vec4<f32> [@location(1)] -> %b1 {
+ %b1 = block {
+ ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
+ }
+}
+)");
+}
+
+TEST_F(IR_BuilderImplTest, EmitFunction_ReturnLocation_Interpolate) {
+ Func("test", utils::Empty, ty.vec4<f32>(), utils::Vector{Return(vec4<f32>(1_f, 2_f, 3_f, 4_f))},
+ utils::Vector{Stage(ast::PipelineStage::kFragment)},
+ utils::Vector{Location(1_i), Interpolate(builtin::InterpolationType::kLinear,
+ builtin::InterpolationSampling::kCentroid)});
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ EXPECT_EQ(
+ Disassemble(m.Get()),
+ R"(%test = @fragment func():vec4<f32> [@location(1), @interpolate(linear, centroid)] -> %b1 {
%b1 = block {
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
}
@@ -145,7 +164,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():f32 [@fragment ra: @frag_depth] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test = @fragment func():f32 [@frag_depth] -> %b1 {
%b1 = block {
ret 1.0f
}
@@ -161,7 +180,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():u32 [@fragment ra: @sample_mask] -> %b1 {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%test = @fragment func():u32 [@sample_mask] -> %b1 {
%b1 = block {
ret 1u
}
diff --git a/src/tint/ir/from_program_store_test.cc b/src/tint/ir/from_program_store_test.cc
index 38efea0..20ea290 100644
--- a/src/tint/ir/from_program_store_test.cc
+++ b/src/tint/ir/from_program_store_test.cc
@@ -40,7 +40,7 @@
%a:ptr<private, u32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
store %a, 4u
ret
diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc
index f9be668..8f5348b 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -144,7 +144,7 @@
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
if true [t: %b2, f: %b3, m: %b4]
# True block
@@ -184,7 +184,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
if true [t: %b2, f: %b3, m: %b4]
# True block
@@ -224,7 +224,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
if true [t: %b2, f: %b3, m: %b4]
# True block
@@ -264,7 +264,7 @@
EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
if true [t: %b2, f: %b3]
# True block
@@ -298,7 +298,7 @@
ASSERT_NE(loop_flow, nullptr);
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
if true [t: %b2, f: %b3, m: %b4]
# True block
@@ -347,7 +347,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, m: %b3]
%b2 = block {
@@ -387,7 +387,7 @@
EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -442,7 +442,7 @@
EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -475,7 +475,7 @@
auto m = res.Move();
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -519,7 +519,7 @@
EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3]
%b2 = block {
@@ -568,7 +568,7 @@
EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2]
%b2 = block {
@@ -606,7 +606,7 @@
EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2]
%b2 = block {
@@ -640,7 +640,7 @@
EXPECT_EQ(0u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, m: %b3]
%b2 = block {
@@ -685,7 +685,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -813,7 +813,7 @@
EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -874,7 +874,7 @@
EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, c: %b3, m: %b4]
%b2 = block {
@@ -967,7 +967,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
loop [s: %b2, m: %b3]
%b2 = block {
@@ -1021,7 +1021,7 @@
EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
switch 1i [c: (0i, %b2), c: (1i, %b3), c: (default, %b4), m: %b5]
# Case block
@@ -1082,7 +1082,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
switch 1i [c: (0i 1i default, %b2), m: %b3]
# Case block
@@ -1121,7 +1121,7 @@
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
switch 1i [c: (default, %b2), m: %b3]
# Case block
@@ -1169,7 +1169,7 @@
// This is 1 because the if is dead-code eliminated and the return doesn't happen.
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
switch 1i [c: (0i, %b2), c: (default, %b3), m: %b4]
# Case block
@@ -1224,7 +1224,7 @@
EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
switch 1i [c: (0i, %b2), c: (default, %b3)]
# Case block
@@ -1255,7 +1255,7 @@
ret 1i
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:i32 = call %b
ret
@@ -1264,5 +1264,85 @@
)");
}
+TEST_F(IR_BuilderImplTest, Func_WithParam_WithAttribute_Invariant) {
+ Func(
+ "f",
+ utils::Vector{Param("a", ty.vec4<f32>(),
+ utils::Vector{Invariant(), Builtin(builtin::BuiltinValue::kPosition)})},
+ ty.vec4<f32>(), utils::Vector{Return("a")},
+ utils::Vector{Stage(ast::PipelineStage::kFragment)}, utils::Vector{Location(1_i)});
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ EXPECT_EQ(
+ Disassemble(m.Get()),
+ R"(%f = @fragment func(%a:vec4<f32> [@invariant, @position]):vec4<f32> [@location(1)] -> %b1 {
+ %b1 = block {
+ ret %a
+ }
+}
+)");
+}
+
+TEST_F(IR_BuilderImplTest, Func_WithParam_WithAttribute_Location) {
+ Func("f", utils::Vector{Param("a", ty.f32(), utils::Vector{Location(2_i)})}, ty.f32(),
+ utils::Vector{Return("a")}, utils::Vector{Stage(ast::PipelineStage::kFragment)},
+ utils::Vector{Location(1_i)});
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%f = @fragment func(%a:f32 [@location(2)]):f32 [@location(1)] -> %b1 {
+ %b1 = block {
+ ret %a
+ }
+}
+)");
+}
+
+TEST_F(IR_BuilderImplTest, Func_WithParam_WithAttribute_Location_WithInterpolation_LinearCentroid) {
+ Func("f",
+ utils::Vector{Param(
+ "a", ty.f32(),
+ utils::Vector{Location(2_i), Interpolate(builtin::InterpolationType::kLinear,
+ builtin::InterpolationSampling::kCentroid)})},
+ ty.f32(), utils::Vector{Return("a")}, utils::Vector{Stage(ast::PipelineStage::kFragment)},
+ utils::Vector{Location(1_i)});
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ EXPECT_EQ(
+ Disassemble(m.Get()),
+ R"(%f = @fragment func(%a:f32 [@location(2), @interpolate(linear, centroid)]):f32 [@location(1)] -> %b1 {
+ %b1 = block {
+ ret %a
+ }
+}
+)");
+}
+
+TEST_F(IR_BuilderImplTest, Func_WithParam_WithAttribute_Location_WithInterpolation_Flat) {
+ Func("f",
+ utils::Vector{
+ Param("a", ty.f32(),
+ utils::Vector{Location(2_i), Interpolate(builtin::InterpolationType::kFlat)})},
+ ty.f32(), utils::Vector{Return("a")}, utils::Vector{Stage(ast::PipelineStage::kFragment)},
+ utils::Vector{Location(1_i)});
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ EXPECT_EQ(
+ Disassemble(m.Get()),
+ R"(%f = @fragment func(%a:f32 [@location(2), @interpolate(flat)]):f32 [@location(1)] -> %b1 {
+ %b1 = block {
+ ret %a
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::ir
diff --git a/src/tint/ir/from_program_unary_test.cc b/src/tint/ir/from_program_unary_test.cc
index bb58c02..b83e5d5 100644
--- a/src/tint/ir/from_program_unary_test.cc
+++ b/src/tint/ir/from_program_unary_test.cc
@@ -39,7 +39,7 @@
ret false
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = call %my_func
%tint_symbol:bool = eq %3, false
@@ -62,7 +62,7 @@
ret 1u
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:u32 = call %my_func
%tint_symbol:u32 = complement %3
@@ -85,7 +85,7 @@
ret 1i
}
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:i32 = call %my_func
%tint_symbol:i32 = negation %3
@@ -109,7 +109,7 @@
%v2:ptr<private, i32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
ret
}
@@ -133,7 +133,7 @@
%v3:ptr<private, i32, read_write> = var
}
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
store %v3, 42i
ret
diff --git a/src/tint/ir/from_program_var_test.cc b/src/tint/ir/from_program_var_test.cc
index c533738..91b8b8f 100644
--- a/src/tint/ir/from_program_var_test.cc
+++ b/src/tint/ir/from_program_var_test.cc
@@ -55,6 +55,21 @@
)");
}
+TEST_F(IR_BuilderImplTest, Emit_GlobalVar_GroupBinding) {
+ GlobalVar("a", ty.u32(), builtin::AddressSpace::kStorage,
+ utils::Vector{Group(2_u), Binding(3_u)});
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
+%b1 = block {
+ %a:ptr<storage, u32, read> = var @binding_point(2, 3)
+}
+
+)");
+}
+
TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
auto* a = Var("a", ty.u32(), builtin::AddressSpace::kFunction);
WrapInFunction(a);
@@ -63,7 +78,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, u32, read_write> = var
ret
@@ -81,7 +96,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, u32, read_write> = var, 2u
ret
@@ -99,7 +114,7 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, u32, read_write> = var
%3:u32 = load %a
diff --git a/src/tint/ir/function.cc b/src/tint/ir/function.cc
index 4e8ea4c..3769d4f 100644
--- a/src/tint/ir/function.cc
+++ b/src/tint/ir/function.cc
@@ -21,7 +21,9 @@
Function::Function(const type::Type* rt,
PipelineStage stage,
std::optional<std::array<uint32_t, 3>> wg_size)
- : Base(), return_type_(rt), pipeline_stage_(stage), workgroup_size_(wg_size) {}
+ : Base(), pipeline_stage_(stage), workgroup_size_(wg_size) {
+ return_.type = rt;
+}
Function::~Function() = default;
@@ -39,20 +41,14 @@
return out << "<unknown>";
}
-utils::StringStream& operator<<(utils::StringStream& out, Function::ReturnAttribute value) {
+utils::StringStream& operator<<(utils::StringStream& out, enum Function::ReturnBuiltin value) {
switch (value) {
- case Function::ReturnAttribute::kLocation:
- return out << "location";
- case Function::ReturnAttribute::kFragDepth:
+ case Function::ReturnBuiltin::kFragDepth:
return out << "frag_depth";
- case Function::ReturnAttribute::kSampleMask:
+ case Function::ReturnBuiltin::kSampleMask:
return out << "sample_mask";
- case Function::ReturnAttribute::kPosition:
+ case Function::ReturnBuiltin::kPosition:
return out << "position";
- case Function::ReturnAttribute::kInvariant:
- return out << "invariant";
- default:
- break;
}
return out << "<unknown>";
}
diff --git a/src/tint/ir/function.h b/src/tint/ir/function.h
index 172afdc..fa5c547 100644
--- a/src/tint/ir/function.h
+++ b/src/tint/ir/function.h
@@ -20,6 +20,7 @@
#include <utility>
#include "src/tint/ir/function_param.h"
+#include "src/tint/ir/location.h"
#include "src/tint/ir/value.h"
#include "src/tint/type/type.h"
@@ -46,20 +47,14 @@
kVertex,
};
- /// Attributes attached to return types
- enum class ReturnAttribute {
- /// No return attribute
- kNone,
- /// Location attribute
- kLocation,
+ /// Builtin attached to return types
+ enum class ReturnBuiltin {
/// Builtin Position attribute
kPosition,
/// Builtin FragDepth attribute
kFragDepth,
/// Builtin SampleMask
kSampleMask,
- /// Invariant attribute
- kInvariant,
};
/// Constructor
@@ -88,26 +83,35 @@
std::optional<std::array<uint32_t, 3>> WorkgroupSize() const { return workgroup_size_; }
/// @returns the return type for the function
- const type::Type* ReturnType() const { return return_type_; }
+ const type::Type* ReturnType() const { return return_.type; }
/// Sets the return attributes
- /// @param attrs the attributes to set
- void SetReturnAttributes(utils::VectorRef<ReturnAttribute> attrs) {
- return_attributes_ = std::move(attrs);
+ /// @param builtin the builtin to set
+ void SetReturnBuiltin(ReturnBuiltin builtin) {
+ TINT_ASSERT(IR, !return_.builtin.has_value());
+ return_.builtin = builtin;
}
- /// @returns the return attributes
- utils::VectorRef<ReturnAttribute> ReturnAttributes() const { return return_attributes_; }
+ /// @returns the return builtin attribute
+ std::optional<enum ReturnBuiltin> ReturnBuiltin() const { return return_.builtin; }
/// Sets the return location
/// @param loc the location to set
- void SetReturnLocation(std::optional<uint32_t> loc) { return_location_ = loc; }
+ /// @param interp the interpolation
+ void SetReturnLocation(uint32_t loc, std::optional<builtin::Interpolation> interp) {
+ return_.location = {loc, interp};
+ }
/// @returns the return location
- std::optional<uint32_t> ReturnLocation() const { return return_location_; }
+ std::optional<Location> ReturnLocation() const { return return_.location; }
+
+ /// Sets the return as invariant
+ /// @param val the invariant value to set
+ void SetReturnInvariant(bool val) { return_.invariant = val; }
+ /// @returns the return invariant value
+ bool ReturnInvariant() const { return return_.invariant; }
/// Sets the function parameters
/// @param params the function paramters
void SetParams(utils::VectorRef<FunctionParam*> params) { params_ = std::move(params); }
-
/// @returns the function parameters
utils::VectorRef<FunctionParam*> Params() const { return params_; }
@@ -118,20 +122,22 @@
Block* StartTarget() const { return start_target_; }
private:
- const type::Type* return_type_;
PipelineStage pipeline_stage_;
std::optional<std::array<uint32_t, 3>> workgroup_size_;
- utils::Vector<ReturnAttribute, 1> return_attributes_;
- std::optional<uint32_t> return_location_;
+ struct {
+ const type::Type* type = nullptr;
+ std::optional<enum ReturnBuiltin> builtin;
+ std::optional<Location> location;
+ bool invariant = false;
+ } return_;
utils::Vector<FunctionParam*, 1> params_;
-
Block* start_target_ = nullptr;
};
utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStage value);
-utils::StringStream& operator<<(utils::StringStream& out, Function::ReturnAttribute value);
+utils::StringStream& operator<<(utils::StringStream& out, enum Function::ReturnBuiltin value);
} // namespace tint::ir
diff --git a/src/tint/ir/function_param.cc b/src/tint/ir/function_param.cc
index a9570fa..ae94f24 100644
--- a/src/tint/ir/function_param.cc
+++ b/src/tint/ir/function_param.cc
@@ -22,4 +22,43 @@
FunctionParam::~FunctionParam() = default;
+utils::StringStream& operator<<(utils::StringStream& out, enum FunctionParam::Builtin value) {
+ switch (value) {
+ case FunctionParam::Builtin::kVertexIndex:
+ out << "vertex_index";
+ break;
+ case FunctionParam::Builtin::kInstanceIndex:
+ out << "instance_index";
+ break;
+ case FunctionParam::Builtin::kPosition:
+ out << "position";
+ break;
+ case FunctionParam::Builtin::kFrontFacing:
+ out << "front_facing";
+ break;
+ case FunctionParam::Builtin::kLocalInvocationId:
+ out << "local_invocation_id";
+ break;
+ case FunctionParam::Builtin::kLocalInvocationIndex:
+ out << "local_invocation_index";
+ break;
+ case FunctionParam::Builtin::kGlobalInvocationId:
+ out << "global_invocation_id";
+ break;
+ case FunctionParam::Builtin::kWorkgroupId:
+ out << "workgroup_id";
+ break;
+ case FunctionParam::Builtin::kNumWorkgroups:
+ out << "num_workgroups";
+ break;
+ case FunctionParam::Builtin::kSampleIndex:
+ out << "sample_index";
+ break;
+ case FunctionParam::Builtin::kSampleMask:
+ out << "sample_mask";
+ break;
+ }
+ return out;
+}
+
} // namespace tint::ir
diff --git a/src/tint/ir/function_param.h b/src/tint/ir/function_param.h
index 2da0584..45ce4e9 100644
--- a/src/tint/ir/function_param.h
+++ b/src/tint/ir/function_param.h
@@ -15,14 +15,45 @@
#ifndef SRC_TINT_IR_FUNCTION_PARAM_H_
#define SRC_TINT_IR_FUNCTION_PARAM_H_
+#include <utility>
+
+#include "src/tint/ir/binding_point.h"
+#include "src/tint/ir/location.h"
#include "src/tint/ir/value.h"
#include "src/tint/utils/castable.h"
+#include "src/tint/utils/vector.h"
namespace tint::ir {
/// A function parameter in the IR.
class FunctionParam : public utils::Castable<FunctionParam, Value> {
public:
+ /// Builtin attribute
+ enum class Builtin {
+ /// Builtin Vertex index
+ kVertexIndex,
+ /// Builtin Instance index
+ kInstanceIndex,
+ /// Builtin Position
+ kPosition,
+ /// Builtin FrontFacing
+ kFrontFacing,
+ /// Builtin Local invocation id
+ kLocalInvocationId,
+ /// Builtin Local invocation index
+ kLocalInvocationIndex,
+ /// Builtin Global invocation id
+ kGlobalInvocationId,
+ /// Builtin Workgroup id
+ kWorkgroupId,
+ /// Builtin Num workgroups
+ kNumWorkgroups,
+ /// Builtin Sample index
+ kSampleIndex,
+ /// Builtin Sample mask
+ kSampleMask,
+ };
+
/// Constructor
/// @param type the type of the var
explicit FunctionParam(const type::Type* type);
@@ -31,11 +62,47 @@
/// @returns the type of the var
const type::Type* Type() const override { return type_; }
+ /// Sets the builtin information. Note, it is currently an error if the builtin is already set.
+ /// @param val the builtin to set
+ void SetBuiltin(FunctionParam::Builtin val) {
+ TINT_ASSERT(IR, !builtin_.has_value());
+ builtin_ = val;
+ }
+ /// @returns the builtin set for the parameter
+ std::optional<FunctionParam::Builtin> Builtin() const { return builtin_; }
+
+ /// Sets the parameter as invariant
+ /// @param val the value to set for invariant
+ void SetInvariant(bool val) { invariant_ = val; }
+ /// @returns true if parameter is invariant
+ bool Invariant() const { return invariant_; }
+
+ /// Sets the location
+ /// @param loc the location value
+ /// @param interpolation if the location interpolation settings
+ void SetLocation(uint32_t loc, std::optional<builtin::Interpolation> interpolation) {
+ location_ = {loc, interpolation};
+ }
+ /// @returns the location if `Attributes` contains `kLocation`
+ std::optional<struct Location> Location() const { return location_; }
+
+ /// Sets the binding point
+ /// @param group the group
+ /// @param binding the binding
+ void SetBindingPoint(uint32_t group, uint32_t binding) { binding_point_ = {group, binding}; }
+ /// @returns the binding points if `Attributes` contains `kBindingPoint`
+ std::optional<struct BindingPoint> BindingPoint() const { return binding_point_; }
+
private:
- /// The type of the parameter
- const type::Type* type_;
+ const type::Type* type_ = nullptr;
+ std::optional<enum FunctionParam::Builtin> builtin_;
+ std::optional<struct Location> location_;
+ std::optional<struct BindingPoint> binding_point_;
+ bool invariant_ = false;
};
+utils::StringStream& operator<<(utils::StringStream& out, enum FunctionParam::Builtin value);
+
} // namespace tint::ir
#endif // SRC_TINT_IR_FUNCTION_PARAM_H_
diff --git a/src/tint/ir/if.cc b/src/tint/ir/if.cc
index 2069c6d..aa1e884 100644
--- a/src/tint/ir/if.cc
+++ b/src/tint/ir/if.cc
@@ -18,7 +18,7 @@
namespace tint::ir {
-If::If(Value* cond, Block* t, Block* f, Block* m)
+If::If(Value* cond, ir::Block* t, ir::Block* f, ir::Block* m)
: Base(utils::Empty), condition_(cond), true_(t), false_(f), merge_(m) {
TINT_ASSERT(IR, true_);
TINT_ASSERT(IR, false_);
diff --git a/src/tint/ir/if.h b/src/tint/ir/if.h
index ad4db1d..6dbe5fc 100644
--- a/src/tint/ir/if.h
+++ b/src/tint/ir/if.h
@@ -34,7 +34,7 @@
/// @param t the true block
/// @param f the false block
/// @param m the merge block
- explicit If(Value* cond, Block* t, Block* f, Block* m);
+ explicit If(Value* cond, ir::Block* t, ir::Block* f, ir::Block* m);
~If() override;
/// @returns the if condition
@@ -43,25 +43,25 @@
Value* Condition() { return condition_; }
/// @returns the true branch block
- const Block* True() const { return true_; }
+ const ir::Block* True() const { return true_; }
/// @returns the true branch block
- Block* True() { return true_; }
+ ir::Block* True() { return true_; }
/// @returns the false branch block
- const Block* False() const { return false_; }
+ const ir::Block* False() const { return false_; }
/// @returns the false branch block
- Block* False() { return false_; }
+ ir::Block* False() { return false_; }
/// @returns the merge branch block
- const Block* Merge() const { return merge_; }
+ const ir::Block* Merge() const { return merge_; }
/// @returns the merge branch block
- Block* Merge() { return merge_; }
+ ir::Block* Merge() { return merge_; }
private:
Value* condition_ = nullptr;
- Block* true_ = nullptr;
- Block* false_ = nullptr;
- Block* merge_ = nullptr;
+ ir::Block* true_ = nullptr;
+ ir::Block* false_ = nullptr;
+ ir::Block* merge_ = nullptr;
};
} // namespace tint::ir
diff --git a/src/tint/ir/instruction.h b/src/tint/ir/instruction.h
index c3c18cd..13bcc00 100644
--- a/src/tint/ir/instruction.h
+++ b/src/tint/ir/instruction.h
@@ -18,6 +18,11 @@
#include "src/tint/ir/value.h"
#include "src/tint/utils/castable.h"
+// Forward declarations
+namespace tint::ir {
+class Block;
+} // namespace tint::ir
+
namespace tint::ir {
/// An instruction in the IR.
@@ -26,9 +31,22 @@
/// Destructor
~Instruction() override;
+ /// Sets the block that owns this instruction
+ /// @param block the new owner block
+ void SetBlock(ir::Block* block) { block_ = block; }
+
+ /// @returns the block that owns this instruction
+ ir::Block* Block() { return block_; }
+
+ /// @returns the block that owns this instruction
+ const ir::Block* Block() const { return block_; }
+
protected:
/// Constructor
Instruction();
+
+ /// The block that owns this instruction
+ ir::Block* block_ = nullptr;
};
} // namespace tint::ir
diff --git a/src/tint/ir/location.h b/src/tint/ir/location.h
new file mode 100644
index 0000000..5edb882
--- /dev/null
+++ b/src/tint/ir/location.h
@@ -0,0 +1,34 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_IR_LOCATION_H_
+#define SRC_TINT_IR_LOCATION_H_
+
+#include <optional>
+
+#include "src/tint/builtin/interpolation.h"
+
+namespace tint::ir {
+
+/// A function parameter in the IR.
+struct Location {
+ /// The location value
+ uint32_t value = 0;
+ /// The interpolation settings
+ std::optional<builtin::Interpolation> interpolation;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_LOCATION_H_
diff --git a/src/tint/ir/loop.cc b/src/tint/ir/loop.cc
index fe34283..bd697bf 100644
--- a/src/tint/ir/loop.cc
+++ b/src/tint/ir/loop.cc
@@ -18,7 +18,7 @@
namespace tint::ir {
-Loop::Loop(Block* s, Block* c, Block* m)
+Loop::Loop(ir::Block* s, ir::Block* c, ir::Block* m)
: Base(utils::Empty), start_(s), continuing_(c), merge_(m) {
TINT_ASSERT(IR, start_);
TINT_ASSERT(IR, continuing_);
diff --git a/src/tint/ir/loop.h b/src/tint/ir/loop.h
index 954a64a..aadfd30 100644
--- a/src/tint/ir/loop.h
+++ b/src/tint/ir/loop.h
@@ -27,28 +27,28 @@
/// @param s the start block
/// @param c the continuing block
/// @param m the merge block
- Loop(Block* s, Block* c, Block* m);
+ Loop(ir::Block* s, ir::Block* c, ir::Block* m);
~Loop() override;
/// @returns the switch start branch
- const Block* Start() const { return start_; }
+ const ir::Block* Start() const { return start_; }
/// @returns the switch start branch
- Block* Start() { return start_; }
+ ir::Block* Start() { return start_; }
/// @returns the switch continuing branch
- const Block* Continuing() const { return continuing_; }
+ const ir::Block* Continuing() const { return continuing_; }
/// @returns the switch continuing branch
- Block* Continuing() { return continuing_; }
+ ir::Block* Continuing() { return continuing_; }
/// @returns the switch merge branch
- const Block* Merge() const { return merge_; }
+ const ir::Block* Merge() const { return merge_; }
/// @returns the switch merge branch
- Block* Merge() { return merge_; }
+ ir::Block* Merge() { return merge_; }
private:
- Block* start_ = nullptr;
- Block* continuing_ = nullptr;
- Block* merge_ = nullptr;
+ ir::Block* start_ = nullptr;
+ ir::Block* continuing_ = nullptr;
+ ir::Block* merge_ = nullptr;
};
} // namespace tint::ir
diff --git a/src/tint/ir/switch.cc b/src/tint/ir/switch.cc
index 003feb2..3756f45 100644
--- a/src/tint/ir/switch.cc
+++ b/src/tint/ir/switch.cc
@@ -18,7 +18,7 @@
namespace tint::ir {
-Switch::Switch(Value* cond, Block* m) : Base(utils::Empty), condition_(cond), merge_(m) {
+Switch::Switch(Value* cond, ir::Block* m) : Base(utils::Empty), condition_(cond), merge_(m) {
TINT_ASSERT(IR, condition_);
TINT_ASSERT(IR, merge_);
condition_->AddUsage(this);
diff --git a/src/tint/ir/switch.h b/src/tint/ir/switch.h
index 588fee9..d9355f1 100644
--- a/src/tint/ir/switch.h
+++ b/src/tint/ir/switch.h
@@ -39,24 +39,24 @@
/// The case selector for this node
utils::Vector<CaseSelector, 4> selectors;
/// The start block for the case block.
- Block* start = nullptr;
+ ir::Block* start = nullptr;
/// @returns the case start target
- const Block* Start() const { return start; }
+ const ir::Block* Start() const { return start; }
/// @returns the case start target
- Block* Start() { return start; }
+ ir::Block* Start() { return start; }
};
/// Constructor
/// @param cond the condition
/// @param m the merge block
- explicit Switch(Value* cond, Block* m);
+ explicit Switch(Value* cond, ir::Block* m);
~Switch() override;
/// @returns the switch merge branch
- const Block* Merge() const { return merge_; }
+ const ir::Block* Merge() const { return merge_; }
/// @returns the switch merge branch
- Block* Merge() { return merge_; }
+ ir::Block* Merge() { return merge_; }
/// @returns the switch cases
utils::VectorRef<Case> Cases() const { return cases_; }
@@ -70,7 +70,7 @@
private:
Value* condition_ = nullptr;
- Block* merge_ = nullptr;
+ ir::Block* merge_ = nullptr;
utils::Vector<Case, 4> cases_;
};
diff --git a/src/tint/ir/transform/add_empty_entry_point_test.cc b/src/tint/ir/transform/add_empty_entry_point_test.cc
index ca918f2..1da2e7d 100644
--- a/src/tint/ir/transform/add_empty_entry_point_test.cc
+++ b/src/tint/ir/transform/add_empty_entry_point_test.cc
@@ -25,7 +25,7 @@
TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) {
auto* expect = R"(
-%unused_entry_point = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+%unused_entry_point = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
ret
}
@@ -43,7 +43,7 @@
mod.functions.Push(ep);
auto* expect = R"(
-%main = func():void [@fragment] -> %b1 {
+%main = @fragment func():void -> %b1 {
%b1 = block {
ret
}
diff --git a/src/tint/ir/var.h b/src/tint/ir/var.h
index 8561e5c..8048463 100644
--- a/src/tint/ir/var.h
+++ b/src/tint/ir/var.h
@@ -17,8 +17,10 @@
#include "src/tint/builtin/access.h"
#include "src/tint/builtin/address_space.h"
+#include "src/tint/ir/binding_point.h"
#include "src/tint/ir/instruction.h"
#include "src/tint/utils/castable.h"
+#include "src/tint/utils/vector.h"
namespace tint::ir {
@@ -39,9 +41,17 @@
/// @returns the initializer
const Value* Initializer() const { return initializer_; }
+ /// Sets the binding point
+ /// @param group the group
+ /// @param binding the binding
+ void SetBindingPoint(uint32_t group, uint32_t binding) { binding_point_ = {group, binding}; }
+ /// @returns the binding points if `Attributes` contains `kBindingPoint`
+ std::optional<struct BindingPoint> BindingPoint() const { return binding_point_; }
+
private:
const type::Type* type_;
Value* initializer_ = nullptr;
+ std::optional<struct BindingPoint> binding_point_;
};
} // namespace tint::ir
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_loop_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_loop_test.cc
index 8ec7a9c..e16d116 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_loop_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_loop_test.cc
@@ -24,11 +24,11 @@
auto* loop = b.CreateLoop();
- loop->Start()->Instructions().Push(b.Continue(loop));
- loop->Continuing()->Instructions().Push(b.BreakIf(b.Constant(true), loop));
- loop->Merge()->Instructions().Push(b.Return(func));
+ loop->Start()->AddInstruction(b.Continue(loop));
+ loop->Continuing()->AddInstruction(b.BreakIf(b.Constant(true), loop));
+ loop->Merge()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(loop);
+ func->StartTarget()->AddInstruction(loop);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -58,10 +58,10 @@
auto* loop = b.CreateLoop();
- loop->Start()->Instructions().Push(b.ExitLoop(loop));
- loop->Merge()->Instructions().Push(b.Return(func));
+ loop->Start()->AddInstruction(b.ExitLoop(loop));
+ loop->Merge()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(loop);
+ func->StartTarget()->AddInstruction(loop);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -89,15 +89,15 @@
auto* loop = b.CreateLoop();
auto* cond_break = b.CreateIf(b.Constant(true));
- cond_break->True()->Instructions().Push(b.ExitLoop(loop));
- cond_break->False()->Instructions().Push(b.ExitIf(cond_break));
- cond_break->Merge()->Instructions().Push(b.Continue(loop));
+ cond_break->True()->AddInstruction(b.ExitLoop(loop));
+ cond_break->False()->AddInstruction(b.ExitIf(cond_break));
+ cond_break->Merge()->AddInstruction(b.Continue(loop));
- loop->Start()->Instructions().Push(cond_break);
- loop->Continuing()->Instructions().Push(b.NextIteration(loop));
- loop->Merge()->Instructions().Push(b.Return(func));
+ loop->Start()->AddInstruction(cond_break);
+ loop->Continuing()->AddInstruction(b.NextIteration(loop));
+ loop->Merge()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(loop);
+ func->StartTarget()->AddInstruction(loop);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -132,15 +132,15 @@
auto* loop = b.CreateLoop();
auto* cond_break = b.CreateIf(b.Constant(true));
- cond_break->True()->Instructions().Push(b.Continue(loop));
- cond_break->False()->Instructions().Push(b.ExitIf(cond_break));
- cond_break->Merge()->Instructions().Push(b.ExitLoop(loop));
+ cond_break->True()->AddInstruction(b.Continue(loop));
+ cond_break->False()->AddInstruction(b.ExitIf(cond_break));
+ cond_break->Merge()->AddInstruction(b.ExitLoop(loop));
- loop->Start()->Instructions().Push(cond_break);
- loop->Continuing()->Instructions().Push(b.NextIteration(loop));
- loop->Merge()->Instructions().Push(b.Return(func));
+ loop->Start()->AddInstruction(cond_break);
+ loop->Continuing()->AddInstruction(b.NextIteration(loop));
+ loop->Merge()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(loop);
+ func->StartTarget()->AddInstruction(loop);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -176,9 +176,9 @@
auto* loop = b.CreateLoop();
- loop->Start()->Instructions().Push(b.Return(func));
+ loop->Start()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(loop);
+ func->StartTarget()->AddInstruction(loop);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -207,11 +207,11 @@
auto* result = b.Equal(mod.Types().i32(), b.Constant(1_i), b.Constant(2_i));
- loop->Start()->Instructions().Push(result);
- loop->Continuing()->Instructions().Push(b.BreakIf(result, loop));
- loop->Merge()->Instructions().Push(b.Return(func));
+ loop->Start()->AddInstruction(result);
+ loop->Continuing()->AddInstruction(b.BreakIf(result, loop));
+ loop->Merge()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(loop);
+ func->StartTarget()->AddInstruction(loop);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -242,15 +242,15 @@
auto* outer_loop = b.CreateLoop();
auto* inner_loop = b.CreateLoop();
- inner_loop->Start()->Instructions().Push(b.ExitLoop(inner_loop));
- inner_loop->Continuing()->Instructions().Push(b.NextIteration(inner_loop));
- inner_loop->Merge()->Instructions().Push(b.Continue(outer_loop));
+ inner_loop->Start()->AddInstruction(b.ExitLoop(inner_loop));
+ inner_loop->Continuing()->AddInstruction(b.NextIteration(inner_loop));
+ inner_loop->Merge()->AddInstruction(b.Continue(outer_loop));
- outer_loop->Start()->Instructions().Push(inner_loop);
- outer_loop->Continuing()->Instructions().Push(b.BreakIf(b.Constant(true), outer_loop));
- outer_loop->Merge()->Instructions().Push(b.Return(func));
+ outer_loop->Start()->AddInstruction(inner_loop);
+ outer_loop->Continuing()->AddInstruction(b.BreakIf(b.Constant(true), outer_loop));
+ outer_loop->Merge()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(outer_loop);
+ func->StartTarget()->AddInstruction(outer_loop);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -289,15 +289,15 @@
auto* outer_loop = b.CreateLoop();
auto* inner_loop = b.CreateLoop();
- inner_loop->Start()->Instructions().Push(b.Continue(inner_loop));
- inner_loop->Continuing()->Instructions().Push(b.BreakIf(b.Constant(true), inner_loop));
- inner_loop->Merge()->Instructions().Push(b.BreakIf(b.Constant(true), outer_loop));
+ inner_loop->Start()->AddInstruction(b.Continue(inner_loop));
+ inner_loop->Continuing()->AddInstruction(b.BreakIf(b.Constant(true), inner_loop));
+ inner_loop->Merge()->AddInstruction(b.BreakIf(b.Constant(true), outer_loop));
- outer_loop->Start()->Instructions().Push(b.Continue(outer_loop));
- outer_loop->Continuing()->Instructions().Push(inner_loop);
- outer_loop->Merge()->Instructions().Push(b.Return(func));
+ outer_loop->Start()->AddInstruction(b.Continue(outer_loop));
+ outer_loop->Continuing()->AddInstruction(inner_loop);
+ outer_loop->Merge()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(outer_loop);
+ func->StartTarget()->AddInstruction(outer_loop);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_switch_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_switch_test.cc
index 5ef5901..92800a3 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_switch_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_switch_test.cc
@@ -25,11 +25,11 @@
auto* swtch = b.CreateSwitch(b.Constant(42_i));
auto* def_case = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector()});
- def_case->Instructions().Push(b.ExitSwitch(swtch));
+ def_case->AddInstruction(b.ExitSwitch(swtch));
- swtch->Merge()->Instructions().Push(b.Return(func));
+ swtch->Merge()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(swtch);
+ func->StartTarget()->AddInstruction(swtch);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -55,17 +55,17 @@
auto* swtch = b.CreateSwitch(b.Constant(42_i));
auto* case_a = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector{b.Constant(1_i)}});
- case_a->Instructions().Push(b.ExitSwitch(swtch));
+ case_a->AddInstruction(b.ExitSwitch(swtch));
auto* case_b = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector{b.Constant(2_i)}});
- case_b->Instructions().Push(b.ExitSwitch(swtch));
+ case_b->AddInstruction(b.ExitSwitch(swtch));
auto* def_case = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector()});
- def_case->Instructions().Push(b.ExitSwitch(swtch));
+ def_case->AddInstruction(b.ExitSwitch(swtch));
- swtch->Merge()->Instructions().Push(b.Return(func));
+ swtch->Merge()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(swtch);
+ func->StartTarget()->AddInstruction(swtch);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -96,19 +96,19 @@
auto* case_a = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector{b.Constant(1_i)},
ir::Switch::CaseSelector{b.Constant(3_i)}});
- case_a->Instructions().Push(b.ExitSwitch(swtch));
+ case_a->AddInstruction(b.ExitSwitch(swtch));
auto* case_b = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector{b.Constant(2_i)},
ir::Switch::CaseSelector{b.Constant(4_i)}});
- case_b->Instructions().Push(b.ExitSwitch(swtch));
+ case_b->AddInstruction(b.ExitSwitch(swtch));
auto* def_case = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector{b.Constant(5_i)},
ir::Switch::CaseSelector()});
- def_case->Instructions().Push(b.ExitSwitch(swtch));
+ def_case->AddInstruction(b.ExitSwitch(swtch));
- swtch->Merge()->Instructions().Push(b.Return(func));
+ swtch->Merge()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(swtch);
+ func->StartTarget()->AddInstruction(swtch);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -138,15 +138,15 @@
auto* swtch = b.CreateSwitch(b.Constant(42_i));
auto* case_a = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector{b.Constant(1_i)}});
- case_a->Instructions().Push(b.Return(func));
+ case_a->AddInstruction(b.Return(func));
auto* case_b = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector{b.Constant(2_i)}});
- case_b->Instructions().Push(b.Return(func));
+ case_b->AddInstruction(b.Return(func));
auto* def_case = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector()});
- def_case->Instructions().Push(b.Return(func));
+ def_case->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(swtch);
+ func->StartTarget()->AddInstruction(swtch);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -176,19 +176,19 @@
auto* swtch = b.CreateSwitch(b.Constant(42_i));
auto* cond_break = b.CreateIf(b.Constant(true));
- cond_break->True()->Instructions().Push(b.ExitSwitch(swtch));
- cond_break->False()->Instructions().Push(b.ExitIf(cond_break));
- cond_break->Merge()->Instructions().Push(b.Return(func));
+ cond_break->True()->AddInstruction(b.ExitSwitch(swtch));
+ cond_break->False()->AddInstruction(b.ExitIf(cond_break));
+ cond_break->Merge()->AddInstruction(b.Return(func));
auto* case_a = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector{b.Constant(1_i)}});
- case_a->Instructions().Push(cond_break);
+ case_a->AddInstruction(cond_break);
auto* def_case = b.CreateCase(swtch, utils::Vector{ir::Switch::CaseSelector()});
- def_case->Instructions().Push(b.ExitSwitch(swtch));
+ def_case->AddInstruction(b.ExitSwitch(swtch));
- swtch->Merge()->Instructions().Push(b.Return(func));
+ swtch->Merge()->AddInstruction(b.Return(func));
- func->StartTarget()->Instructions().Push(swtch);
+ func->StartTarget()->AddInstruction(swtch);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
