Import Tint changes from Dawn
Changes:
- 9f83fa1d1be3af162b8ef8d35d7edffc5bdc3484 tint/ir: Remove 'type' field from Instruction by Ben Clayton <bclayton@google.com>
- c95576ed7b8609b0baa12bd912d374fa56b496ef [tint] Fix syntax_tree build by Ben Clayton <bclayton@google.com>
- 42de29e08832ac2422fc308cc16e62d87282fd2a tint/ir: Add an initializer field to Var by Ben Clayton <bclayton@google.com>
- 88417687fac18974ee168271f3e8e526c86e4ffa tint/ir: Add name metadata to Module by Ben Clayton <bclayton@google.com>
GitOrigin-RevId: 9f83fa1d1be3af162b8ef8d35d7edffc5bdc3484
Change-Id: Ib0c53763bf910764302c5f7540f60382e922c914
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/132405
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 92edf12..0a6f186 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -2161,6 +2161,7 @@
"ir/builder_impl_var_test.cc",
"ir/constant_test.cc",
"ir/discard_test.cc",
+ "ir/module_test.cc",
"ir/store_test.cc",
"ir/test_helper.h",
"ir/unary_test.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 60b9885..0b06c66 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -1457,6 +1457,7 @@
ir/builder_impl_var_test.cc
ir/constant_test.cc
ir/discard_test.cc
+ ir/module_test.cc
ir/store_test.cc
ir/test_helper.h
ir/unary_test.cc
diff --git a/src/tint/ir/binary.cc b/src/tint/ir/binary.cc
index 063924f..af28306 100644
--- a/src/tint/ir/binary.cc
+++ b/src/tint/ir/binary.cc
@@ -19,10 +19,10 @@
namespace tint::ir {
-Binary::Binary(Kind kind, const type::Type* ty, Value* lhs, Value* rhs)
- : Base(ty), kind_(kind), lhs_(lhs), rhs_(rhs) {
- TINT_ASSERT(IR, lhs_);
- TINT_ASSERT(IR, rhs_);
+Binary::Binary(Kind k, const type::Type* res_ty, Value* lhs, Value* rhs)
+ : kind(k), result_type(res_ty), lhs_(lhs), rhs_(rhs) {
+ TINT_ASSERT(IR, lhs);
+ TINT_ASSERT(IR, rhs);
lhs_->AddUsage(this);
rhs_->AddUsage(this);
}
diff --git a/src/tint/ir/binary.h b/src/tint/ir/binary.h
index e941b3f..a6dc722 100644
--- a/src/tint/ir/binary.h
+++ b/src/tint/ir/binary.h
@@ -59,8 +59,8 @@
Binary& operator=(const Binary& inst) = delete;
Binary& operator=(Binary&& inst) = delete;
- /// @returns the kind of instruction
- Kind GetKind() const { return kind_; }
+ /// @returns the type of the value
+ const type::Type* Type() const override { return result_type; }
/// @returns the left-hand-side value for the instruction
const Value* LHS() const { return lhs_; }
@@ -68,8 +68,13 @@
/// @returns the right-hand-side value for the instruction
const Value* RHS() const { return rhs_; }
+ /// the kind of binary instruction
+ Kind kind = Kind::kAdd;
+
+ /// the result type of the instruction
+ const type::Type* result_type = nullptr;
+
private:
- Kind kind_;
Value* lhs_ = nullptr;
Value* rhs_ = nullptr;
};
diff --git a/src/tint/ir/binary_test.cc b/src/tint/ir/binary_test.cc
index 6a96633..7baf0cd 100644
--- a/src/tint/ir/binary_test.cc
+++ b/src/tint/ir/binary_test.cc
@@ -29,7 +29,8 @@
b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kAnd);
+ EXPECT_EQ(inst->kind, Binary::Kind::kAnd);
+ ASSERT_NE(inst->result_type, nullptr);
ASSERT_NE(inst->Type(), nullptr);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
@@ -50,7 +51,7 @@
b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kOr);
+ EXPECT_EQ(inst->kind, Binary::Kind::kOr);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -70,7 +71,7 @@
b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kXor);
+ EXPECT_EQ(inst->kind, Binary::Kind::kXor);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -90,7 +91,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kEqual);
+ EXPECT_EQ(inst->kind, Binary::Kind::kEqual);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -110,7 +111,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kNotEqual);
+ EXPECT_EQ(inst->kind, Binary::Kind::kNotEqual);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -130,7 +131,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kLessThan);
+ EXPECT_EQ(inst->kind, Binary::Kind::kLessThan);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -150,7 +151,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kGreaterThan);
+ EXPECT_EQ(inst->kind, Binary::Kind::kGreaterThan);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -170,7 +171,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kLessThanEqual);
+ EXPECT_EQ(inst->kind, Binary::Kind::kLessThanEqual);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -190,7 +191,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kGreaterThanEqual);
+ EXPECT_EQ(inst->kind, Binary::Kind::kGreaterThanEqual);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -209,7 +210,7 @@
b.builder.Not(b.builder.ir.types.Get<type::Bool>(), b.builder.Constant(true));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kEqual);
+ EXPECT_EQ(inst->kind, Binary::Kind::kEqual);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -229,7 +230,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kShiftLeft);
+ EXPECT_EQ(inst->kind, Binary::Kind::kShiftLeft);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -249,7 +250,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kShiftRight);
+ EXPECT_EQ(inst->kind, Binary::Kind::kShiftRight);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -269,7 +270,7 @@
b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kAdd);
+ EXPECT_EQ(inst->kind, Binary::Kind::kAdd);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -289,7 +290,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kSubtract);
+ EXPECT_EQ(inst->kind, Binary::Kind::kSubtract);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -309,7 +310,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kMultiply);
+ EXPECT_EQ(inst->kind, Binary::Kind::kMultiply);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -329,7 +330,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kDivide);
+ EXPECT_EQ(inst->kind, Binary::Kind::kDivide);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -349,7 +350,7 @@
b.builder.Constant(4_i), b.builder.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kModulo);
+ EXPECT_EQ(inst->kind, Binary::Kind::kModulo);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
auto lhs = inst->LHS()->As<Constant>()->value;
@@ -367,7 +368,7 @@
const auto* inst = b.builder.And(b.builder.ir.types.Get<type::I32>(), b.builder.Constant(4_i),
b.builder.Constant(2_i));
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kAnd);
+ EXPECT_EQ(inst->kind, Binary::Kind::kAnd);
ASSERT_NE(inst->LHS(), nullptr);
ASSERT_EQ(inst->LHS()->Usage().Length(), 1u);
@@ -383,7 +384,7 @@
auto val = b.builder.Constant(4_i);
const auto* inst = b.builder.And(b.builder.ir.types.Get<type::I32>(), val, val);
- EXPECT_EQ(inst->GetKind(), Binary::Kind::kAnd);
+ EXPECT_EQ(inst->kind, Binary::Kind::kAnd);
ASSERT_EQ(inst->LHS(), inst->RHS());
ASSERT_NE(inst->LHS(), nullptr);
diff --git a/src/tint/ir/builder_impl.cc b/src/tint/ir/builder_impl.cc
index ef4a94e..bf6a0e7 100644
--- a/src/tint/ir/builder_impl.cc
+++ b/src/tint/ir/builder_impl.cc
@@ -775,12 +775,13 @@
if (!init) {
return;
}
-
- auto* store = builder.Store(val, init.Get());
- current_flow_block->instructions.Push(store);
+ val->initializer = init.Get();
}
// Store the declaration so we can get the instruction to store too
scopes_.Set(v->name->symbol, val);
+
+ // Record the original name of the var
+ builder.ir.SetName(val, v->name->symbol.Name());
},
[&](const ast::Let* l) {
// A `let` doesn't exist as a standalone item in the IR, it's just the result of the
@@ -792,6 +793,9 @@
// Store the results of the initialization
scopes_.Set(l->name->symbol, init.Get());
+
+ // Record the original name of the let
+ builder.ir.SetName(init.Get(), l->name->symbol.Name());
},
[&](const ast::Override*) {
add_error(var->source,
diff --git a/src/tint/ir/builder_impl_binary_test.cc b/src/tint/ir/builder_impl_binary_test.cc
index ed13159..3be8405 100644
--- a/src/tint/ir/builder_impl_binary_test.cc
+++ b/src/tint/ir/builder_impl_binary_test.cc
@@ -54,14 +54,14 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ref<private, u32, read_write> = add %1:ref<private, u32, read_write>, 1u
- store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+ %2:ref<private, u32, read_write> = add %v1:ref<private, u32, read_write>, 1u
+ store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
ret
func_end
@@ -96,14 +96,14 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ref<private, u32, read_write> = sub %1:ref<private, u32, read_write>, 1u
- store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+ %2:ref<private, u32, read_write> = sub %v1:ref<private, u32, read_write>, 1u
+ store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
ret
func_end
@@ -138,14 +138,14 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ref<private, u32, read_write> = mul %1:ref<private, u32, read_write>, 1u
- store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+ %2:ref<private, u32, read_write> = mul %v1:ref<private, u32, read_write>, 1u
+ store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
ret
func_end
@@ -180,14 +180,14 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ref<private, u32, read_write> = div %1:ref<private, u32, read_write>, 1u
- store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+ %2:ref<private, u32, read_write> = div %v1:ref<private, u32, read_write>, 1u
+ store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
ret
func_end
@@ -222,14 +222,14 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ref<private, u32, read_write> = mod %1:ref<private, u32, read_write>, 1u
- store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+ %2:ref<private, u32, read_write> = mod %v1:ref<private, u32, read_write>, 1u
+ store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
ret
func_end
@@ -264,14 +264,14 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, bool, read_write> = var private read_write
+%v1:ref<private, bool, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ref<private, bool, read_write> = and %1:ref<private, bool, read_write>, false
- store %1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
+ %2:ref<private, bool, read_write> = and %v1:ref<private, bool, read_write>, false
+ store %v1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
ret
func_end
@@ -306,14 +306,14 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, bool, read_write> = var private read_write
+%v1:ref<private, bool, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ref<private, bool, read_write> = or %1:ref<private, bool, read_write>, false
- store %1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
+ %2:ref<private, bool, read_write> = or %v1:ref<private, bool, read_write>, false
+ store %v1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
ret
func_end
@@ -348,14 +348,14 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ref<private, u32, read_write> = xor %1:ref<private, u32, read_write>, 1u
- store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+ %2:ref<private, u32, read_write> = xor %v1:ref<private, u32, read_write>, 1u
+ store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
ret
func_end
@@ -379,14 +379,14 @@
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn4 = block
%1:bool = call my_func
- %2:bool = var function read_write
- store %2:bool, %1:bool
+ %tint_symbol:bool = var function, read_write
+ store %tint_symbol:bool, %1:bool
branch %fn5
%fn5 = if %1:bool [t: %fn6, f: %fn7, m: %fn8]
# true branch
%fn6 = block
- store %2:bool, false
+ store %tint_symbol:bool, false
branch %fn8
# if merge
@@ -414,15 +414,15 @@
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn4 = block
%1:bool = call my_func
- %2:bool = var function read_write
- store %2:bool, %1:bool
+ %tint_symbol:bool = var function, read_write
+ store %tint_symbol:bool, %1:bool
branch %fn5
%fn5 = if %1:bool [t: %fn6, f: %fn7, m: %fn8]
# true branch
# false branch
%fn7 = block
- store %2:bool, true
+ store %tint_symbol:bool, true
branch %fn8
# if merge
@@ -569,14 +569,14 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ref<private, u32, read_write> = shiftl %1:ref<private, u32, read_write>, 1u
- store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+ %2:ref<private, u32, read_write> = shiftl %v1:ref<private, u32, read_write>, 1u
+ store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
ret
func_end
@@ -611,14 +611,14 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, u32, read_write> = var private read_write
+%v1:ref<private, u32, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ref<private, u32, read_write> = shiftr %1:ref<private, u32, read_write>, 1u
- store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
+ %2:ref<private, u32, read_write> = shiftr %v1:ref<private, u32, read_write>, 1u
+ store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
ret
func_end
@@ -645,8 +645,8 @@
%fn4 = block
%1:f32 = call my_func
%2:bool = lt %1:f32, 2.0f
- %3:bool = var function read_write
- store %3:bool, %2:bool
+ %tint_symbol:bool = var function, read_write
+ store %tint_symbol:bool, %2:bool
branch %fn5
%fn5 = if %2:bool [t: %fn6, f: %fn7, m: %fn8]
@@ -657,7 +657,7 @@
%6:f32 = mul 2.29999995231628417969f, %5:f32
%7:f32 = div %4:f32, %6:f32
%8:bool = gt 2.5f, %7:f32
- store %3:bool, %8:bool
+ store %tint_symbol:bool, %8:bool
branch %fn8
# if merge
@@ -685,7 +685,7 @@
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn4 = block
- %1:bool = call my_func, false
+ %tint_symbol:bool = call my_func, false
ret
func_end
diff --git a/src/tint/ir/builder_impl_call_test.cc b/src/tint/ir/builder_impl_call_test.cc
index 6339626..0fb070e 100644
--- a/src/tint/ir/builder_impl_call_test.cc
+++ b/src/tint/ir/builder_impl_call_test.cc
@@ -92,14 +92,13 @@
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, i32, read_write> = var private read_write
-store %1:ref<private, i32, read_write>, 1i
+%i:ref<private, i32, read_write> = var private, read_write, 1i
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:f32 = convert i32, %1:ref<private, i32, read_write>
+ %tint_symbol:f32 = convert i32, %i:ref<private, i32, read_write>
ret
func_end
@@ -116,8 +115,7 @@
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, vec3<f32>, read_write> = var private read_write
-store %1:ref<private, vec3<f32>, read_write>, vec3<f32> 0.0f
+%i:ref<private, vec3<f32>, read_write> = var private, read_write, vec3<f32> 0.0f
@@ -135,14 +133,13 @@
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, f32, read_write> = var private read_write
-store %1:ref<private, f32, read_write>, 1.0f
+%i:ref<private, f32, read_write> = var private, read_write, 1.0f
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:vec3<f32> = construct 2.0f, 3.0f, %1:ref<private, f32, read_write>
+ %tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %i:ref<private, f32, read_write>
ret
func_end
diff --git a/src/tint/ir/builder_impl_store_test.cc b/src/tint/ir/builder_impl_store_test.cc
index 7cd074c..82fbc0a 100644
--- a/src/tint/ir/builder_impl_store_test.cc
+++ b/src/tint/ir/builder_impl_store_test.cc
@@ -37,13 +37,13 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, u32, read_write> = var private read_write
+%a:ref<private, u32, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- store %1:ref<private, u32, read_write>, 4u
+ store %a:ref<private, u32, read_write>, 4u
ret
func_end
diff --git a/src/tint/ir/builder_impl_unary_test.cc b/src/tint/ir/builder_impl_unary_test.cc
index 091493d..2323acb 100644
--- a/src/tint/ir/builder_impl_unary_test.cc
+++ b/src/tint/ir/builder_impl_unary_test.cc
@@ -91,13 +91,13 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, i32, read_write> = var private read_write
+%v1:ref<private, i32, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ptr<private, i32, read_write> = addr_of %1:ref<private, i32, read_write>
+ %v2:ptr<private, i32, read_write> = addr_of %v1:ref<private, i32, read_write>
ret
func_end
@@ -117,14 +117,14 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, i32, read_write> = var private read_write
+%v1:ref<private, i32, read_write> = var private, read_write
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
- %2:ptr<private, i32, read_write> = addr_of %1:ref<private, i32, read_write>
- %3:i32 = indirection %2:ptr<private, i32, read_write>
+ %v3:ptr<private, i32, read_write> = addr_of %v1:ref<private, i32, read_write>
+ %v2:i32 = indirection %v3:ptr<private, i32, read_write>
ret
func_end
diff --git a/src/tint/ir/builder_impl_var_test.cc b/src/tint/ir/builder_impl_var_test.cc
index d68ce8a..a3f4d27 100644
--- a/src/tint/ir/builder_impl_var_test.cc
+++ b/src/tint/ir/builder_impl_var_test.cc
@@ -34,7 +34,7 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, u32, read_write> = var private read_write
+%a:ref<private, u32, read_write> = var private, read_write
@@ -50,8 +50,7 @@
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
-%1:ref<private, u32, read_write> = var private read_write
-store %1:ref<private, u32, read_write>, 2u
+%a:ref<private, u32, read_write> = var private, read_write, 2u
@@ -69,7 +68,7 @@
EXPECT_EQ(Disassemble(m),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn2 = block
- %1:ref<function, u32, read_write> = var function read_write
+ %a:ref<function, u32, read_write> = var function, read_write
ret
func_end
@@ -88,8 +87,7 @@
EXPECT_EQ(Disassemble(m),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn2 = block
- %1:ref<function, u32, read_write> = var function read_write
- store %1:ref<function, u32, read_write>, 2u
+ %a:ref<function, u32, read_write> = var function, read_write, 2u
ret
func_end
diff --git a/src/tint/ir/builtin.h b/src/tint/ir/builtin.h
index 9a82dba..de0e435 100644
--- a/src/tint/ir/builtin.h
+++ b/src/tint/ir/builtin.h
@@ -25,10 +25,10 @@
class Builtin : public utils::Castable<Builtin, Call> {
public:
/// Constructor
- /// @param type the result type
+ /// @param res_type the result type
/// @param func the builtin function
/// @param args the conversion arguments
- Builtin(const type::Type* type, builtin::Function func, utils::VectorRef<Value*> args);
+ Builtin(const type::Type* res_type, builtin::Function func, utils::VectorRef<Value*> args);
Builtin(const Builtin& inst) = delete;
Builtin(Builtin&& inst) = delete;
~Builtin() override;
diff --git a/src/tint/ir/call.cc b/src/tint/ir/call.cc
index d8c1af6..dd503ea 100644
--- a/src/tint/ir/call.cc
+++ b/src/tint/ir/call.cc
@@ -20,8 +20,8 @@
namespace tint::ir {
-Call::Call(const type::Type* ty, utils::VectorRef<Value*> arguments)
- : Base(ty), args(std::move(arguments)) {
+Call::Call(const type::Type* res_ty, utils::VectorRef<Value*> arguments)
+ : result_type(res_ty), args(std::move(arguments)) {
for (auto* arg : args) {
arg->AddUsage(this);
}
diff --git a/src/tint/ir/call.h b/src/tint/ir/call.h
index fde193e..5393810 100644
--- a/src/tint/ir/call.h
+++ b/src/tint/ir/call.h
@@ -30,6 +30,12 @@
Call& operator=(const Call& inst) = delete;
Call& operator=(Call&& inst) = delete;
+ /// @returns the type of the value
+ const type::Type* Type() const override { return result_type; }
+
+ /// The instruction type
+ const type::Type* result_type = nullptr;
+
/// The constructor arguments
utils::Vector<Value*, 1> args;
@@ -37,9 +43,9 @@
/// Constructor
Call() = delete;
/// Constructor
- /// @param type the result type
+ /// @param result_type the result type
/// @param args the constructor arguments
- Call(const type::Type* type, utils::VectorRef<Value*> args);
+ Call(const type::Type* result_type, utils::VectorRef<Value*> args);
};
} // namespace tint::ir
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 50d11cc..688fc41 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -94,7 +94,12 @@
std::string_view Disassembler::IdOf(const Value* value) {
TINT_ASSERT(IR, value);
- return value_ids_.GetOrCreate(value, [&] { return std::to_string(value_ids_.Count()); });
+ return value_ids_.GetOrCreate(value, [&] {
+ if (auto sym = mod_.NameOf(value)) {
+ return sym.Name();
+ }
+ return std::to_string(value_ids_.Count());
+ });
}
void Disassembler::Walk(const FlowNode* node) {
@@ -401,7 +406,11 @@
},
[&](const ir::Var* v) {
EmitValue(v);
- out_ << " = var " << v->address_space << " " << v->access;
+ out_ << " = var " << v->address_space << ", " << v->access;
+ if (v->initializer) {
+ out_ << ", ";
+ EmitValue(v->initializer);
+ }
});
}
@@ -419,7 +428,7 @@
void Disassembler::EmitBinary(const Binary* b) {
EmitValue(b);
out_ << " = ";
- switch (b->GetKind()) {
+ switch (b->kind) {
case Binary::Kind::kAdd:
out_ << "add";
break;
@@ -478,7 +487,7 @@
void Disassembler::EmitUnary(const Unary* u) {
EmitValue(u);
out_ << " = ";
- switch (u->GetKind()) {
+ switch (u->kind) {
case Unary::Kind::kAddressOf:
out_ << "addr_of";
break;
diff --git a/src/tint/ir/instruction.cc b/src/tint/ir/instruction.cc
index bbd4992..e54b13f 100644
--- a/src/tint/ir/instruction.cc
+++ b/src/tint/ir/instruction.cc
@@ -20,8 +20,6 @@
Instruction::Instruction() = default;
-Instruction::Instruction(const type::Type* ty) : type(ty) {}
-
Instruction::~Instruction() = default;
} // namespace tint::ir
diff --git a/src/tint/ir/instruction.h b/src/tint/ir/instruction.h
index 1bfde50..8b52f01 100644
--- a/src/tint/ir/instruction.h
+++ b/src/tint/ir/instruction.h
@@ -31,18 +31,9 @@
Instruction& operator=(const Instruction& inst) = delete;
Instruction& operator=(Instruction&& inst) = delete;
- /// @returns the type of the value
- const type::Type* Type() const override { return type; }
-
- /// The instruction type
- const type::Type* type = nullptr;
-
protected:
/// Constructor
Instruction();
- /// Constructor
- /// @param type the result type
- explicit Instruction(const type::Type* type);
};
} // namespace tint::ir
diff --git a/src/tint/ir/module.cc b/src/tint/ir/module.cc
index 133d364..acee4ae 100644
--- a/src/tint/ir/module.cc
+++ b/src/tint/ir/module.cc
@@ -14,6 +14,8 @@
#include "src/tint/ir/module.h"
+#include <limits>
+
namespace tint::ir {
Module::Module() = default;
@@ -24,4 +26,33 @@
Module& Module::operator=(Module&&) = default;
+Symbol Module::NameOf(const Value* value) const {
+ return value_to_id_.Get(value).value_or(Symbol{});
+}
+
+Symbol Module::SetName(const Value* value, std::string_view name) {
+ TINT_ASSERT(IR, !name.empty());
+
+ if (auto old = value_to_id_.Get(value)) {
+ value_to_id_.Remove(value);
+ id_to_value_.Remove(old.value());
+ }
+
+ auto sym = symbols.Register(name);
+ if (id_to_value_.Add(sym, value)) {
+ value_to_id_.Add(value, sym);
+ return sym;
+ }
+ auto prefix = std::string(name) + "_";
+ for (uint64_t suffix = 1; suffix != std::numeric_limits<uint64_t>::max(); suffix++) {
+ sym = symbols.Register(prefix + std::to_string(suffix));
+ if (id_to_value_.Add(sym, value)) {
+ value_to_id_.Add(value, sym);
+ return sym;
+ }
+ }
+ TINT_ASSERT(IR, false); // !
+ return Symbol{};
+}
+
} // namespace tint::ir
diff --git a/src/tint/ir/module.h b/src/tint/ir/module.h
index ebf9209..0f89d86 100644
--- a/src/tint/ir/module.h
+++ b/src/tint/ir/module.h
@@ -32,6 +32,15 @@
/// Main module class for the IR.
class Module {
+ /// Program Id required to create other components
+ ProgramID prog_id_;
+
+ /// Map of value to pre-declared identifier
+ utils::Hashmap<const Value*, Symbol, 32> value_to_id_;
+
+ /// Map of pre-declared identifier to value
+ utils::Hashmap<Symbol, const Value*, 32> id_to_value_;
+
public:
/// Constructor
Module();
@@ -46,11 +55,15 @@
/// @returns a reference to this module
Module& operator=(Module&& o);
- private:
- /// Program Id required to create other components
- ProgramID prog_id_;
+ /// @param value the value
+ /// @return the name of the given value, or an invalid symbol if the value is not named.
+ Symbol NameOf(const Value* value) const;
- public:
+ /// @param value the value to name.
+ /// @param name the desired name of the value. May be suffixed on collision.
+ /// @return the unique symbol of the given value.
+ Symbol SetName(const Value* value, std::string_view name);
+
/// The flow node allocator
utils::BlockAllocator<FlowNode> flow_nodes;
/// The constant allocator
diff --git a/src/tint/ir/module_test.cc b/src/tint/ir/module_test.cc
new file mode 100644
index 0000000..15d5cb1
--- /dev/null
+++ b/src/tint/ir/module_test.cc
@@ -0,0 +1,66 @@
+// 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.
+
+#include "src/tint/ir/module.h"
+#include "src/tint/ir/test_helper.h"
+
+namespace tint::ir {
+namespace {
+
+using namespace tint::number_suffixes; // NOLINT
+
+using IR_ModuleTest = TestHelper;
+
+TEST_F(IR_ModuleTest, NameOfUnnamed) {
+ Module mod;
+ auto* v = mod.values.Create<ir::Var>(
+ mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+ EXPECT_FALSE(mod.NameOf(v).IsValid());
+}
+
+TEST_F(IR_ModuleTest, SetName) {
+ Module mod;
+ auto* v = mod.values.Create<ir::Var>(
+ mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+ EXPECT_EQ(mod.SetName(v, "a").Name(), "a");
+ EXPECT_EQ(mod.NameOf(v).Name(), "a");
+}
+
+TEST_F(IR_ModuleTest, SetNameRename) {
+ Module mod;
+ auto* v = mod.values.Create<ir::Var>(
+ mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+ EXPECT_EQ(mod.SetName(v, "a").Name(), "a");
+ EXPECT_EQ(mod.SetName(v, "b").Name(), "b");
+ EXPECT_EQ(mod.NameOf(v).Name(), "b");
+}
+
+TEST_F(IR_ModuleTest, SetNameCollision) {
+ Module mod;
+ auto* a = mod.values.Create<ir::Var>(
+ mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+ auto* b = mod.values.Create<ir::Var>(
+ mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+ auto* c = mod.values.Create<ir::Var>(
+ mod.types.Get<type::I32>(), builtin::AddressSpace::kUndefined, builtin::Access::kUndefined);
+ EXPECT_EQ(mod.SetName(a, "x").Name(), "x");
+ EXPECT_EQ(mod.SetName(b, "x_1").Name(), "x_1");
+ EXPECT_EQ(mod.SetName(c, "x").Name(), "x_2");
+ EXPECT_EQ(mod.NameOf(a).Name(), "x");
+ EXPECT_EQ(mod.NameOf(b).Name(), "x_1");
+ EXPECT_EQ(mod.NameOf(c).Name(), "x_2");
+}
+
+} // namespace
+} // namespace tint::ir
diff --git a/src/tint/ir/unary.cc b/src/tint/ir/unary.cc
index 2fd7e57..db13d5d 100644
--- a/src/tint/ir/unary.cc
+++ b/src/tint/ir/unary.cc
@@ -19,7 +19,8 @@
namespace tint::ir {
-Unary::Unary(Kind kind, const type::Type* ty, Value* val) : Base(ty), kind_(kind), val_(val) {
+Unary::Unary(Kind k, const type::Type* res_ty, Value* val)
+ : kind(k), result_type(res_ty), val_(val) {
TINT_ASSERT(IR, val_);
val_->AddUsage(this);
}
diff --git a/src/tint/ir/unary.h b/src/tint/ir/unary.h
index 8eb0bb4..e665386 100644
--- a/src/tint/ir/unary.h
+++ b/src/tint/ir/unary.h
@@ -33,9 +33,9 @@
/// Constructor
/// @param kind the kind of unary instruction
- /// @param type the result type
- /// @param val the lhs of the instruction
- Unary(Kind kind, const type::Type* type, Value* val);
+ /// @param result_type the result type
+ /// @param val the input value for the instruction
+ Unary(Kind kind, const type::Type* result_type, Value* val);
Unary(const Unary& inst) = delete;
Unary(Unary&& inst) = delete;
~Unary() override;
@@ -43,14 +43,19 @@
Unary& operator=(const Unary& inst) = delete;
Unary& operator=(Unary&& inst) = delete;
- /// @returns the kind of instruction
- Kind GetKind() const { return kind_; }
+ /// @returns the type of the value
+ const type::Type* Type() const override { return result_type; }
/// @returns the value for the instruction
const Value* Val() const { return val_; }
+ /// the kind of unary instruction
+ Kind kind = Kind::kAddressOf;
+
+ /// the result type of the instruction
+ const type::Type* result_type = nullptr;
+
private:
- Kind kind_;
Value* val_ = nullptr;
};
diff --git a/src/tint/ir/unary_test.cc b/src/tint/ir/unary_test.cc
index e205d8f..86a3993 100644
--- a/src/tint/ir/unary_test.cc
+++ b/src/tint/ir/unary_test.cc
@@ -33,7 +33,7 @@
b.builder.Constant(4_i));
ASSERT_TRUE(inst->Is<Unary>());
- EXPECT_EQ(inst->GetKind(), Unary::Kind::kAddressOf);
+ EXPECT_EQ(inst->kind, Unary::Kind::kAddressOf);
ASSERT_NE(inst->Type(), nullptr);
@@ -49,7 +49,7 @@
b.builder.Complement(b.builder.ir.types.Get<type::I32>(), b.builder.Constant(4_i));
ASSERT_TRUE(inst->Is<Unary>());
- EXPECT_EQ(inst->GetKind(), Unary::Kind::kComplement);
+ EXPECT_EQ(inst->kind, Unary::Kind::kComplement);
ASSERT_TRUE(inst->Val()->Is<Constant>());
auto lhs = inst->Val()->As<Constant>()->value;
@@ -65,7 +65,7 @@
b.builder.Indirection(b.builder.ir.types.Get<type::I32>(), b.builder.Constant(4_i));
ASSERT_TRUE(inst->Is<Unary>());
- EXPECT_EQ(inst->GetKind(), Unary::Kind::kIndirection);
+ EXPECT_EQ(inst->kind, Unary::Kind::kIndirection);
ASSERT_TRUE(inst->Val()->Is<Constant>());
auto lhs = inst->Val()->As<Constant>()->value;
@@ -79,7 +79,7 @@
b.builder.Negation(b.builder.ir.types.Get<type::I32>(), b.builder.Constant(4_i));
ASSERT_TRUE(inst->Is<Unary>());
- EXPECT_EQ(inst->GetKind(), Unary::Kind::kNegation);
+ EXPECT_EQ(inst->kind, Unary::Kind::kNegation);
ASSERT_TRUE(inst->Val()->Is<Constant>());
auto lhs = inst->Val()->As<Constant>()->value;
@@ -92,7 +92,7 @@
const auto* inst =
b.builder.Negation(b.builder.ir.types.Get<type::I32>(), b.builder.Constant(4_i));
- EXPECT_EQ(inst->GetKind(), Unary::Kind::kNegation);
+ EXPECT_EQ(inst->kind, Unary::Kind::kNegation);
ASSERT_NE(inst->Val(), nullptr);
ASSERT_EQ(inst->Val()->Usage().Length(), 1u);
diff --git a/src/tint/ir/value.h b/src/tint/ir/value.h
index e91e4c4..4b7810b 100644
--- a/src/tint/ir/value.h
+++ b/src/tint/ir/value.h
@@ -47,7 +47,7 @@
utils::VectorRef<const Instruction*> Usage() const { return uses_; }
/// @returns the type of the value
- virtual const type::Type* Type() const = 0;
+ virtual const type::Type* Type() const { return nullptr; }
protected:
/// Constructor
diff --git a/src/tint/ir/var.cc b/src/tint/ir/var.cc
index 3ab97e0..e33da54 100644
--- a/src/tint/ir/var.cc
+++ b/src/tint/ir/var.cc
@@ -20,7 +20,7 @@
namespace tint::ir {
Var::Var(const type::Type* ty, builtin::AddressSpace addr_space, builtin::Access acc)
- : Base(ty), address_space(addr_space), access(acc) {}
+ : type(ty), address_space(addr_space), access(acc) {}
Var::~Var() = default;
diff --git a/src/tint/ir/var.h b/src/tint/ir/var.h
index 1387c60..5a61104 100644
--- a/src/tint/ir/var.h
+++ b/src/tint/ir/var.h
@@ -26,7 +26,7 @@
class Var : public utils::Castable<Var, Instruction> {
public:
/// Constructor
- /// @param type the type
+ /// @param type the type of the var
/// @param address_space the address space of the var
/// @param access the access mode of the var
Var(const type::Type* type, builtin::AddressSpace address_space, builtin::Access access);
@@ -37,11 +37,20 @@
Var& operator=(const Var& inst) = delete;
Var& operator=(Var&& inst) = delete;
+ /// @returns the type of the var
+ const type::Type* Type() const override { return type; }
+
+ /// the result type of the instruction
+ const type::Type* type = nullptr;
+
/// The variable address space
builtin::AddressSpace address_space = builtin::AddressSpace::kUndefined;
/// The variable access mode
builtin::Access access = builtin::Access::kUndefined;
+
+ /// The optional initializer
+ Value* initializer = nullptr;
};
} // namespace tint::ir
diff --git a/src/tint/symbol.h b/src/tint/symbol.h
index ae29559..ed6fe72 100644
--- a/src/tint/symbol.h
+++ b/src/tint/symbol.h
@@ -100,6 +100,9 @@
/// @returns true if the symbol is valid
bool IsValid() const { return val_ != static_cast<uint32_t>(-1); }
+ /// @returns true if the symbol is valid
+ operator bool() const { return IsValid(); }
+
/// @returns the value for the symbol
uint32_t value() const { return val_; }
diff --git a/src/tint/writer/spirv/generator_impl_ir.cc b/src/tint/writer/spirv/generator_impl_ir.cc
index d0d7cc0..7ff56be 100644
--- a/src/tint/writer/spirv/generator_impl_ir.cc
+++ b/src/tint/writer/spirv/generator_impl_ir.cc
@@ -241,14 +241,14 @@
// Determine the opcode.
spv::Op op = spv::Op::Max;
- switch (binary->GetKind()) {
+ switch (binary->kind) {
case ir::Binary::Kind::kAdd: {
op = binary->Type()->is_integer_scalar_or_vector() ? spv::Op::OpIAdd : spv::Op::OpFAdd;
break;
}
default: {
TINT_ICE(Writer, diagnostics_)
- << "unimplemented binary instruction: " << static_cast<uint32_t>(binary->GetKind());
+ << "unimplemented binary instruction: " << static_cast<uint32_t>(binary->kind);
}
}
diff --git a/src/tint/writer/syntax_tree/generator_impl.cc b/src/tint/writer/syntax_tree/generator_impl.cc
index 08ca8c5..56f7f57 100644
--- a/src/tint/writer/syntax_tree/generator_impl.cc
+++ b/src/tint/writer/syntax_tree/generator_impl.cc
@@ -976,14 +976,14 @@
case 0: // No initializer
break;
case 1: // Single line initializer statement
- line() << TrimSuffix(init_buf.lines[0].content, ";");
+ line() << utils::TrimSuffix(init_buf.lines[0].content, ";");
break;
default: // Block initializer statement
for (size_t i = 1; i < init_buf.lines.size(); i++) {
// Indent all by the first line
init_buf.lines[i].indent += current_buffer_->current_indent;
}
- line() << TrimSuffix(init_buf.String(), "\n");
+ line() << utils::TrimSuffix(init_buf.String(), "\n");
break;
}
}
@@ -1004,14 +1004,14 @@
case 0: // No continuing
break;
case 1: // Single line continuing statement
- line() << TrimSuffix(cont_buf.lines[0].content, ";");
+ line() << utils::TrimSuffix(cont_buf.lines[0].content, ";");
break;
default: // Block continuing statement
for (size_t i = 1; i < cont_buf.lines.size(); i++) {
// Indent all by the first line
cont_buf.lines[i].indent += current_buffer_->current_indent;
}
- line() << TrimSuffix(cont_buf.String(), "\n");
+ line() << utils::TrimSuffix(cont_buf.String(), "\n");
break;
}
}
