[ir] Remove references, indirection and address-of
Change the type of a `var` to a pointer. Fold away address-of and indirection. Fixed: tint:1912 Change-Id: Ib1f07538c63df9349c5b6171062b6f79750c1439 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/133400 Commit-Queue: James Price <jrprice@google.com> Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
This commit is contained in:
parent
7511bf6334
commit
0bb1bb3067
|
@ -185,18 +185,10 @@ Unary* Builder::CreateUnary(Unary::Kind kind, const type::Type* type, Value* val
|
||||||
return ir.values.Create<ir::Unary>(kind, type, val);
|
return ir.values.Create<ir::Unary>(kind, type, val);
|
||||||
}
|
}
|
||||||
|
|
||||||
Unary* Builder::AddressOf(const type::Type* type, Value* val) {
|
|
||||||
return CreateUnary(Unary::Kind::kAddressOf, type, val);
|
|
||||||
}
|
|
||||||
|
|
||||||
Unary* Builder::Complement(const type::Type* type, Value* val) {
|
Unary* Builder::Complement(const type::Type* type, Value* val) {
|
||||||
return CreateUnary(Unary::Kind::kComplement, type, val);
|
return CreateUnary(Unary::Kind::kComplement, type, val);
|
||||||
}
|
}
|
||||||
|
|
||||||
Unary* Builder::Indirection(const type::Type* type, Value* val) {
|
|
||||||
return CreateUnary(Unary::Kind::kIndirection, type, val);
|
|
||||||
}
|
|
||||||
|
|
||||||
Unary* Builder::Negation(const type::Type* type, Value* val) {
|
Unary* Builder::Negation(const type::Type* type, Value* val) {
|
||||||
return CreateUnary(Unary::Kind::kNegation, type, val);
|
return CreateUnary(Unary::Kind::kNegation, type, val);
|
||||||
}
|
}
|
||||||
|
|
|
@ -279,24 +279,12 @@ class Builder {
|
||||||
/// @returns the operation
|
/// @returns the operation
|
||||||
Unary* CreateUnary(Unary::Kind kind, const type::Type* type, Value* val);
|
Unary* CreateUnary(Unary::Kind kind, const type::Type* type, Value* val);
|
||||||
|
|
||||||
/// Creates an AddressOf operation
|
|
||||||
/// @param type the result type of the expression
|
|
||||||
/// @param val the value
|
|
||||||
/// @returns the operation
|
|
||||||
Unary* AddressOf(const type::Type* type, Value* val);
|
|
||||||
|
|
||||||
/// Creates a Complement operation
|
/// Creates a Complement operation
|
||||||
/// @param type the result type of the expression
|
/// @param type the result type of the expression
|
||||||
/// @param val the value
|
/// @param val the value
|
||||||
/// @returns the operation
|
/// @returns the operation
|
||||||
Unary* Complement(const type::Type* type, Value* val);
|
Unary* Complement(const type::Type* type, Value* val);
|
||||||
|
|
||||||
/// Creates an Indirection operation
|
|
||||||
/// @param type the result type of the expression
|
|
||||||
/// @param val the value
|
|
||||||
/// @returns the operation
|
|
||||||
Unary* Indirection(const type::Type* type, Value* val);
|
|
||||||
|
|
||||||
/// Creates a Negation operation
|
/// Creates a Negation operation
|
||||||
/// @param type the result type of the expression
|
/// @param type the result type of the expression
|
||||||
/// @param val the value
|
/// @param val the value
|
||||||
|
|
|
@ -530,15 +530,9 @@ void Disassembler::EmitUnary(const Unary* u) {
|
||||||
EmitValue(u);
|
EmitValue(u);
|
||||||
out_ << " = ";
|
out_ << " = ";
|
||||||
switch (u->kind) {
|
switch (u->kind) {
|
||||||
case Unary::Kind::kAddressOf:
|
|
||||||
out_ << "addr_of";
|
|
||||||
break;
|
|
||||||
case Unary::Kind::kComplement:
|
case Unary::Kind::kComplement:
|
||||||
out_ << "complement";
|
out_ << "complement";
|
||||||
break;
|
break;
|
||||||
case Unary::Kind::kIndirection:
|
|
||||||
out_ << "indirection";
|
|
||||||
break;
|
|
||||||
case Unary::Kind::kNegation:
|
case Unary::Kind::kNegation:
|
||||||
out_ << "negation";
|
out_ << "negation";
|
||||||
break;
|
break;
|
||||||
|
|
|
@ -82,6 +82,7 @@
|
||||||
#include "src/tint/sem/value_expression.h"
|
#include "src/tint/sem/value_expression.h"
|
||||||
#include "src/tint/sem/variable.h"
|
#include "src/tint/sem/variable.h"
|
||||||
#include "src/tint/switch.h"
|
#include "src/tint/switch.h"
|
||||||
|
#include "src/tint/type/pointer.h"
|
||||||
#include "src/tint/type/reference.h"
|
#include "src/tint/type/reference.h"
|
||||||
#include "src/tint/type/void.h"
|
#include "src/tint/type/void.h"
|
||||||
#include "src/tint/utils/defer.h"
|
#include "src/tint/utils/defer.h"
|
||||||
|
@ -853,7 +854,11 @@ class Impl {
|
||||||
return tint::Switch( //
|
return tint::Switch( //
|
||||||
var,
|
var,
|
||||||
[&](const ast::Var* v) {
|
[&](const ast::Var* v) {
|
||||||
auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx);
|
auto* ref = sem->Type()->As<type::Reference>();
|
||||||
|
auto* ty = builder_.ir.types.Get<type::Pointer>(
|
||||||
|
ref->StoreType()->Clone(clone_ctx_.type_ctx), ref->AddressSpace(),
|
||||||
|
ref->Access());
|
||||||
|
|
||||||
auto* val = builder_.Declare(ty);
|
auto* val = builder_.Declare(ty);
|
||||||
current_flow_block_->instructions.Push(val);
|
current_flow_block_->instructions.Push(val);
|
||||||
|
|
||||||
|
@ -915,14 +920,12 @@ class Impl {
|
||||||
Instruction* inst = nullptr;
|
Instruction* inst = nullptr;
|
||||||
switch (expr->op) {
|
switch (expr->op) {
|
||||||
case ast::UnaryOp::kAddressOf:
|
case ast::UnaryOp::kAddressOf:
|
||||||
inst = builder_.AddressOf(ty, val.Get());
|
case ast::UnaryOp::kIndirection:
|
||||||
break;
|
// 'address-of' and 'indirection' just fold away and we propagate the pointer.
|
||||||
|
return val;
|
||||||
case ast::UnaryOp::kComplement:
|
case ast::UnaryOp::kComplement:
|
||||||
inst = builder_.Complement(ty, val.Get());
|
inst = builder_.Complement(ty, val.Get());
|
||||||
break;
|
break;
|
||||||
case ast::UnaryOp::kIndirection:
|
|
||||||
inst = builder_.Indirection(ty, val.Get());
|
|
||||||
break;
|
|
||||||
case ast::UnaryOp::kNegation:
|
case ast::UnaryOp::kNegation:
|
||||||
inst = builder_.Negation(ty, val.Get());
|
inst = builder_.Negation(ty, val.Get());
|
||||||
break;
|
break;
|
||||||
|
|
|
@ -58,14 +58,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Increment) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, u32, read_write> = var
|
%v1:ptr<private, u32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, u32, read_write> = add %v1:ref<private, u32, read_write>, 1u
|
%2:ptr<private, u32, read_write> = add %v1:ptr<private, u32, read_write>, 1u
|
||||||
store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -81,14 +81,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAdd) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, u32, read_write> = var
|
%v1:ptr<private, u32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, u32, read_write> = add %v1:ref<private, u32, read_write>, 1u
|
%2:ptr<private, u32, read_write> = add %v1:ptr<private, u32, read_write>, 1u
|
||||||
store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -127,14 +127,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Decrement) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, i32, read_write> = var
|
%v1:ptr<private, i32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, i32, read_write> = sub %v1:ref<private, i32, read_write>, 1i
|
%2:ptr<private, i32, read_write> = sub %v1:ptr<private, i32, read_write>, 1u
|
||||||
store %v1:ref<private, i32, read_write>, %2:ref<private, i32, read_write>
|
store %v1:ptr<private, i32, read_write>, %2:ptr<private, i32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -150,14 +150,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundSubtract) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, u32, read_write> = var
|
%v1:ptr<private, u32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, u32, read_write> = sub %v1:ref<private, u32, read_write>, 1u
|
%2:ptr<private, u32, read_write> = sub %v1:ptr<private, u32, read_write>, 1u
|
||||||
store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -196,14 +196,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundMultiply) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, u32, read_write> = var
|
%v1:ptr<private, u32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, u32, read_write> = mul %v1:ref<private, u32, read_write>, 1u
|
%2:ptr<private, u32, read_write> = mul %v1:ptr<private, u32, read_write>, 1u
|
||||||
store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -242,14 +242,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundDiv) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, u32, read_write> = var
|
%v1:ptr<private, u32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, u32, read_write> = div %v1:ref<private, u32, read_write>, 1u
|
%2:ptr<private, u32, read_write> = div %v1:ptr<private, u32, read_write>, 1u
|
||||||
store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -288,14 +288,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundModulo) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, u32, read_write> = var
|
%v1:ptr<private, u32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, u32, read_write> = mod %v1:ref<private, u32, read_write>, 1u
|
%2:ptr<private, u32, read_write> = mod %v1:ptr<private, u32, read_write>, 1u
|
||||||
store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -334,14 +334,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAnd) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, bool, read_write> = var
|
%v1:ptr<private, bool, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, bool, read_write> = and %v1:ref<private, bool, read_write>, false
|
%2:ptr<private, bool, read_write> = and %v1:ptr<private, bool, read_write>, false
|
||||||
store %v1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
|
store %v1:ptr<private, bool, read_write>, %2:ptr<private, bool, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -380,14 +380,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundOr) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, bool, read_write> = var
|
%v1:ptr<private, bool, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, bool, read_write> = or %v1:ref<private, bool, read_write>, false
|
%2:ptr<private, bool, read_write> = or %v1:ptr<private, bool, read_write>, false
|
||||||
store %v1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
|
store %v1:ptr<private, bool, read_write>, %2:ptr<private, bool, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -426,14 +426,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundXor) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, u32, read_write> = var
|
%v1:ptr<private, u32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, u32, read_write> = xor %v1:ref<private, u32, read_write>, 1u
|
%2:ptr<private, u32, read_write> = xor %v1:ptr<private, u32, read_write>, 1u
|
||||||
store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -706,14 +706,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftLeft) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, u32, read_write> = var
|
%v1:ptr<private, u32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, u32, read_write> = shiftl %v1:ref<private, u32, read_write>, 1u
|
%2:ptr<private, u32, read_write> = shiftl %v1:ptr<private, u32, read_write>, 1u
|
||||||
store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -752,14 +752,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftRight) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, u32, read_write> = var
|
%v1:ptr<private, u32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%2:ref<private, u32, read_write> = shiftr %v1:ref<private, u32, read_write>, 1u
|
%2:ptr<private, u32, read_write> = shiftr %v1:ptr<private, u32, read_write>, 1u
|
||||||
store %v1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
store %v1:ptr<private, u32, read_write>, %2:ptr<private, u32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
|
|
@ -100,13 +100,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%i:ref<private, i32, read_write> = var, 1i
|
%i:ptr<private, i32, read_write> = var, 1i
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%tint_symbol:f32 = convert i32, %i:ref<private, i32, read_write>
|
%tint_symbol:f32 = convert i32, %i:ptr<private, i32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -121,7 +121,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%i:ref<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f
|
%i:ptr<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
@ -137,13 +137,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%i:ref<private, f32, read_write> = var, 1.0f
|
%i:ptr<private, f32, read_write> = var, 1.0f
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %i:ref<private, f32, read_write>
|
%tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %i:ptr<private, f32, read_write>
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
|
|
@ -36,13 +36,13 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%a:ref<private, u32, read_write> = var
|
%a:ptr<private, u32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
store %a:ref<private, u32, read_write>, 4u
|
store %a:ptr<private, u32, read_write>, 4u
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
|
|
@ -105,13 +105,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, i32, read_write> = var
|
%v2:ptr<private, i32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%v2:ptr<private, i32, read_write> = addr_of %v1:ref<private, i32, read_write>
|
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -122,7 +121,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
|
||||||
GlobalVar("v1", builtin::AddressSpace::kPrivate, ty.i32());
|
GlobalVar("v1", builtin::AddressSpace::kPrivate, ty.i32());
|
||||||
utils::Vector stmts = {
|
utils::Vector stmts = {
|
||||||
Decl(Let("v3", AddressOf("v1"))),
|
Decl(Let("v3", AddressOf("v1"))),
|
||||||
Decl(Let("v2", Deref("v3"))),
|
Assign(Deref("v3"), 42_i),
|
||||||
};
|
};
|
||||||
WrapInFunction(stmts);
|
WrapInFunction(stmts);
|
||||||
|
|
||||||
|
@ -130,14 +129,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%v1:ref<private, i32, read_write> = var
|
%v3:ptr<private, i32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn3 = block {
|
%fn3 = block {
|
||||||
%v3:ptr<private, i32, read_write> = addr_of %v1:ref<private, i32, read_write>
|
store %v3:ptr<private, i32, read_write>, 42i
|
||||||
%v2:i32 = indirection %v3:ptr<private, i32, read_write>
|
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
|
|
@ -33,7 +33,7 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%a:ref<private, u32, read_write> = var
|
%a:ptr<private, u32, read_write> = var
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
@ -48,7 +48,7 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) {
|
||||||
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
|
||||||
|
|
||||||
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
|
||||||
%a:ref<private, u32, read_write> = var, 2u
|
%a:ptr<private, u32, read_write> = var, 2u
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|
||||||
|
@ -65,7 +65,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
|
||||||
EXPECT_EQ(Disassemble(m.Get()),
|
EXPECT_EQ(Disassemble(m.Get()),
|
||||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn2 = block {
|
%fn2 = block {
|
||||||
%a:ref<function, u32, read_write> = var
|
%a:ptr<function, u32, read_write> = var
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
@ -83,7 +83,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
|
||||||
EXPECT_EQ(Disassemble(m.Get()),
|
EXPECT_EQ(Disassemble(m.Get()),
|
||||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
|
||||||
%fn2 = block {
|
%fn2 = block {
|
||||||
%a:ref<function, u32, read_write> = var, 2u
|
%a:ptr<function, u32, read_write> = var, 2u
|
||||||
} -> %func_end # return
|
} -> %func_end # return
|
||||||
} %func_end
|
} %func_end
|
||||||
|
|
||||||
|
|
|
@ -272,12 +272,12 @@ class State {
|
||||||
|
|
||||||
const ast::VariableDeclStatement* Var(const ir::Var* var) {
|
const ast::VariableDeclStatement* Var(const ir::Var* var) {
|
||||||
Symbol name = NameOf(var);
|
Symbol name = NameOf(var);
|
||||||
auto* ptr = var->Type()->As<type::Reference>();
|
auto* ptr = var->Type()->As<type::Pointer>();
|
||||||
if (!ptr) {
|
if (!ptr) {
|
||||||
Err("Incorrect type for var");
|
Err("Incorrect type for var");
|
||||||
return nullptr;
|
return nullptr;
|
||||||
}
|
}
|
||||||
auto ty = Type(ptr);
|
auto ty = Type(ptr->StoreType());
|
||||||
const ast::Expression* init = nullptr;
|
const ast::Expression* init = nullptr;
|
||||||
if (var->initializer) {
|
if (var->initializer) {
|
||||||
init = Expr(var->initializer);
|
init = Expr(var->initializer);
|
||||||
|
@ -432,7 +432,10 @@ class State {
|
||||||
: builtin::Access::kUndefined;
|
: builtin::Access::kUndefined;
|
||||||
return b.ty.pointer(el.Get(), address_space, access);
|
return b.ty.pointer(el.Get(), address_space, access);
|
||||||
},
|
},
|
||||||
[&](const type::Reference* r) { return Type(r->StoreType()); },
|
[&](const type::Reference*) -> utils::Result<ast::Type> {
|
||||||
|
TINT_ICE(IR, b.Diagnostics()) << "reference types should never appear in the IR";
|
||||||
|
return ast::Type{};
|
||||||
|
},
|
||||||
[&](Default) {
|
[&](Default) {
|
||||||
UNHANDLED_CASE(ty);
|
UNHANDLED_CASE(ty);
|
||||||
return ast::Type{};
|
return ast::Type{};
|
||||||
|
|
|
@ -25,9 +25,7 @@ class Unary : public utils::Castable<Unary, Instruction> {
|
||||||
public:
|
public:
|
||||||
/// The kind of instruction.
|
/// The kind of instruction.
|
||||||
enum class Kind {
|
enum class Kind {
|
||||||
kAddressOf,
|
|
||||||
kComplement,
|
kComplement,
|
||||||
kIndirection,
|
|
||||||
kNegation,
|
kNegation,
|
||||||
};
|
};
|
||||||
|
|
||||||
|
@ -50,7 +48,7 @@ class Unary : public utils::Castable<Unary, Instruction> {
|
||||||
const Value* Val() const { return val_; }
|
const Value* Val() const { return val_; }
|
||||||
|
|
||||||
/// the kind of unary instruction
|
/// the kind of unary instruction
|
||||||
Kind kind = Kind::kAddressOf;
|
Kind kind = Kind::kNegation;
|
||||||
|
|
||||||
/// the result type of the instruction
|
/// the result type of the instruction
|
||||||
const type::Type* result_type = nullptr;
|
const type::Type* result_type = nullptr;
|
||||||
|
|
|
@ -23,27 +23,6 @@ using namespace tint::number_suffixes; // NOLINT
|
||||||
|
|
||||||
using IR_InstructionTest = TestHelper;
|
using IR_InstructionTest = TestHelper;
|
||||||
|
|
||||||
TEST_F(IR_InstructionTest, CreateAddressOf) {
|
|
||||||
Module mod;
|
|
||||||
Builder b{mod};
|
|
||||||
|
|
||||||
// TODO(dsinclair): This would be better as an identifier, but works for now.
|
|
||||||
const auto* inst = b.AddressOf(
|
|
||||||
b.ir.types.Get<type::Pointer>(b.ir.types.Get<type::I32>(), builtin::AddressSpace::kPrivate,
|
|
||||||
builtin::Access::kReadWrite),
|
|
||||||
b.Constant(4_i));
|
|
||||||
|
|
||||||
ASSERT_TRUE(inst->Is<Unary>());
|
|
||||||
EXPECT_EQ(inst->kind, Unary::Kind::kAddressOf);
|
|
||||||
|
|
||||||
ASSERT_NE(inst->Type(), nullptr);
|
|
||||||
|
|
||||||
ASSERT_TRUE(inst->Val()->Is<Constant>());
|
|
||||||
auto lhs = inst->Val()->As<Constant>()->value;
|
|
||||||
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
|
|
||||||
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
|
|
||||||
}
|
|
||||||
|
|
||||||
TEST_F(IR_InstructionTest, CreateComplement) {
|
TEST_F(IR_InstructionTest, CreateComplement) {
|
||||||
Module mod;
|
Module mod;
|
||||||
Builder b{mod};
|
Builder b{mod};
|
||||||
|
@ -58,22 +37,6 @@ TEST_F(IR_InstructionTest, CreateComplement) {
|
||||||
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
|
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
|
||||||
}
|
}
|
||||||
|
|
||||||
TEST_F(IR_InstructionTest, CreateIndirection) {
|
|
||||||
Module mod;
|
|
||||||
Builder b{mod};
|
|
||||||
|
|
||||||
// TODO(dsinclair): This would be better as an identifier, but works for now.
|
|
||||||
const auto* inst = b.Indirection(b.ir.types.Get<type::I32>(), b.Constant(4_i));
|
|
||||||
|
|
||||||
ASSERT_TRUE(inst->Is<Unary>());
|
|
||||||
EXPECT_EQ(inst->kind, Unary::Kind::kIndirection);
|
|
||||||
|
|
||||||
ASSERT_TRUE(inst->Val()->Is<Constant>());
|
|
||||||
auto lhs = inst->Val()->As<Constant>()->value;
|
|
||||||
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
|
|
||||||
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
|
|
||||||
}
|
|
||||||
|
|
||||||
TEST_F(IR_InstructionTest, CreateNegation) {
|
TEST_F(IR_InstructionTest, CreateNegation) {
|
||||||
Module mod;
|
Module mod;
|
||||||
Builder b{mod};
|
Builder b{mod};
|
||||||
|
|
Loading…
Reference in New Issue