[ir] Add multi-element swizzle.
This CL adds support for multi-element swizzle.
Bug: tint:1913
Change-Id: I72dc831fcda041e764b0d1a509ed031d5e66849f
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/135680
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 8200833..8434ede 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1271,6 +1271,8 @@
"ir/store.h",
"ir/switch.cc",
"ir/switch.h",
+ "ir/swizzle.cc",
+ "ir/swizzle.h",
"ir/transform/transform.cc",
"ir/transform/transform.h",
"ir/unary.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 7b13895..586b903 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -780,6 +780,8 @@
ir/store.h
ir/switch.cc
ir/switch.h
+ ir/swizzle.cc
+ ir/swizzle.h
ir/to_program.cc
ir/to_program.h
ir/unary.cc
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index bf42d5d..b93cb94 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -247,4 +247,10 @@
return ir.values.Create<ir::Access>(type, source, indices);
}
+ir::Swizzle* Builder::Swizzle(const type::Type* type,
+ Value* source,
+ utils::VectorRef<uint32_t> indices) {
+ return ir.values.Create<ir::Swizzle>(type, source, indices);
+}
+
} // namespace tint::ir
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index 6afac25..ae20e42 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -42,6 +42,7 @@
#include "src/tint/ir/return.h"
#include "src/tint/ir/store.h"
#include "src/tint/ir/switch.h"
+#include "src/tint/ir/swizzle.h"
#include "src/tint/ir/unary.h"
#include "src/tint/ir/user_call.h"
#include "src/tint/ir/value.h"
@@ -396,6 +397,13 @@
/// @returns the instruction
ir::Access* Access(const type::Type* type, Value* source, utils::VectorRef<Value*> indices);
+ /// Creates a new `Swizzle`
+ /// @param type the return type
+ /// @param source the source value
+ /// @param indices the access indices
+ /// @returns the instruction
+ ir::Swizzle* Swizzle(const type::Type* type, Value* source, utils::VectorRef<uint32_t> indices);
+
/// Retrieves the root block for the module, creating if necessary
/// @returns the root block
ir::Block* CreateRootBlockIfNeeded();
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 7763366..40fbb2e 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -38,6 +38,7 @@
#include "src/tint/ir/return.h"
#include "src/tint/ir/store.h"
#include "src/tint/ir/switch.h"
+#include "src/tint/ir/swizzle.h"
#include "src/tint/ir/user_call.h"
#include "src/tint/ir/var.h"
#include "src/tint/switch.h"
@@ -400,6 +401,29 @@
}
out_ << std::endl;
},
+ [&](const ir::Swizzle* s) {
+ EmitValueWithType(s);
+ out_ << " = swizzle ";
+ EmitValue(s->Object());
+ out_ << ", ";
+ for (auto idx : s->Indices()) {
+ switch (idx) {
+ case 0:
+ out_ << "x";
+ break;
+ case 1:
+ out_ << "y";
+ break;
+ case 2:
+ out_ << "z";
+ break;
+ case 3:
+ out_ << "w";
+ break;
+ }
+ }
+ out_ << std::endl;
+ },
[&](const ir::Branch* b) { EmitBranch(b); },
[&](Default) { out_ << "Unknown instruction: " << inst->TypeInfo().name; });
}
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index 54e78fc..33b888b 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -878,6 +878,8 @@
struct AccessorInfo {
Value* object = nullptr;
+ Instruction* result = nullptr;
+ const type::Type* result_type = nullptr;
utils::Vector<Value*, 1> indices;
};
@@ -905,6 +907,13 @@
info.object = res.Get();
}
+ if (auto* sem = program_->Sem().Get(expr)->As<sem::Load>()) {
+ auto* ref = sem->ReferenceType();
+ info.result_type = ref->StoreType()->Clone(clone_ctx_.type_ctx);
+ } else {
+ info.result_type = program_->Sem().Get(expr)->Type()->Clone(clone_ctx_.type_ctx);
+ }
+
// The AST chain is `inside-out` compared to what we need, which means the list it generates
// is backwards. We need to operate on the list in reverse order to have the correct access
// chain.
@@ -912,10 +921,10 @@
bool ok = tint::Switch(
accessor,
[&](const ast::IndexAccessorExpression* idx) {
- return GenerateIndexAccessor(idx, &info);
+ return GenerateIndexAccessor(idx, info);
},
[&](const ast::MemberAccessorExpression* member) {
- return GenerateMemberAccessor(member, &info);
+ return GenerateMemberAccessor(member, info);
},
[&](Default) {
TINT_ICE(Writer, diagnostics_)
@@ -927,49 +936,85 @@
}
}
- const type::Type* ty = nullptr;
- if (auto* sem = program_->Sem().Get(expr)->As<sem::Load>()) {
- auto* ref = sem->ReferenceType();
- ty = builder_.ir.Types().pointer(ref->StoreType()->Clone(clone_ctx_.type_ctx),
- ref->AddressSpace(), ref->Access());
- } else {
- ty = program_->Sem().Get(expr)->UnwrapLoad()->Type()->Clone(clone_ctx_.type_ctx);
+ if (!info.indices.IsEmpty()) {
+ info.result = GenerateAccess(info);
}
-
- auto* access = builder_.Access(ty, info.object, info.indices);
- current_block_->Append(access);
- return access;
+ return info.result;
}
- bool GenerateIndexAccessor(const ast::IndexAccessorExpression* expr, AccessorInfo* info) {
+ Instruction* GenerateAccess(const AccessorInfo& info) {
+ // The access result type should match the source result type. If the source is a pointer,
+ // we generate a pointer.
+ const type::Type* ty = nullptr;
+ if (info.object->Type()->Is<type::Pointer>() && !info.result_type->Is<type::Pointer>()) {
+ auto* ptr = info.object->Type()->As<type::Pointer>();
+ ty = builder_.ir.Types().pointer(info.result_type, ptr->AddressSpace(), ptr->Access());
+ } else {
+ ty = info.result_type;
+ }
+
+ auto* a = builder_.Access(ty, info.object, info.indices);
+ current_block_->Append(a);
+ return a;
+ }
+
+ bool GenerateIndexAccessor(const ast::IndexAccessorExpression* expr, AccessorInfo& info) {
auto res = EmitExpression(expr->index);
if (!res) {
return false;
}
- info->indices.Push(res.Get());
+ info.indices.Push(res.Get());
return true;
}
- bool GenerateMemberAccessor(const ast::MemberAccessorExpression* expr, AccessorInfo* info) {
+ bool GenerateMemberAccessor(const ast::MemberAccessorExpression* expr, AccessorInfo& info) {
auto* expr_sem = program_->Sem().Get(expr)->UnwrapLoad();
return tint::Switch(
expr_sem, //
[&](const sem::StructMemberAccess* access) {
uint32_t idx = access->Member()->Index();
- info->indices.Push(builder_.Constant(u32(idx)));
+ info.indices.Push(builder_.Constant(u32(idx)));
return true;
},
[&](const sem::Swizzle* swizzle) {
auto& indices = swizzle->Indices();
+
+ // A single element swizzle is just treated as an accessor.
if (indices.Length() == 1) {
- info->indices.Push(builder_.Constant(u32(indices[0])));
+ info.indices.Push(builder_.Constant(u32(indices[0])));
return true;
}
- TINT_ICE(IR, diagnostics_) << "unhandled multi index swizzle";
- return false;
+ // Store the result type away, this will be the result of the swizzle, but the
+ // intermediate steps need different result types.
+ auto* result_type = info.result_type;
+
+ // Emit any preceeding member/index accessors
+ if (!info.indices.IsEmpty()) {
+ // The access chain is being split, the initial part of than will have a
+ // resulting type that matches the object being swizzled.
+ info.result_type = swizzle->Object()->Type()->Clone(clone_ctx_.type_ctx);
+ info.object = GenerateAccess(info);
+ info.indices.Clear();
+
+ // If the sub-accessor generated a pointer result, make sure a load is emitted
+ if (auto* ptr = info.object->Type()->As<type::Pointer>()) {
+ auto* load = builder_.Load(info.object);
+ info.result_type = ptr->StoreType();
+ info.object = load;
+ current_block_->Append(load);
+ }
+ }
+
+ info.result = builder_.Swizzle(swizzle->Type()->Clone(clone_ctx_.type_ctx),
+ info.object, std::move(indices));
+ current_block_->Append(info.result);
+
+ info.object = info.result;
+ info.result_type = result_type;
+ return true;
},
[&](Default) {
TINT_ICE(IR, diagnostics_)
diff --git a/src/tint/ir/from_program_accessor_test.cc b/src/tint/ir/from_program_accessor_test.cc
index 1130c13..4318e94 100644
--- a/src/tint/ir/from_program_accessor_test.cc
+++ b/src/tint/ir/from_program_accessor_test.cc
@@ -193,7 +193,7 @@
)");
}
-TEST_F(IR_FromProgramAccessorTest, DISABLED_Accessor_Var_MultiElementSwizzle) {
+TEST_F(IR_FromProgramAccessorTest, Accessor_Var_MultiElementSwizzle) {
// var a: vec3<f32>
// let b = a.zyxz
@@ -216,7 +216,7 @@
)");
}
-TEST_F(IR_FromProgramAccessorTest, DISABLED_Accessor_Var_MultiElementSwizzleOfSwizzle) {
+TEST_F(IR_FromProgramAccessorTest, Accessor_Var_MultiElementSwizzleOfSwizzle) {
// var a: vec3<f32>
// let b = a.zyx.yy
@@ -232,15 +232,15 @@
%b1 = block {
%a:ptr<function, vec3<f32>, read_write> = var
%3:vec3<f32> = load %a
- %2:vec3<f32> = swizzle %1, zyx
- %b:vec2<f32> = swizzle %2, yy
+ %4:vec3<f32> = swizzle %3, zyx
+ %b:vec2<f32> = swizzle %4, yy
ret
}
}
)");
}
-TEST_F(IR_FromProgramAccessorTest, DISABLED_Accessor_Var_MultiElementSwizzle_MiddleOfChain) {
+TEST_F(IR_FromProgramAccessorTest, Accessor_Var_MultiElementSwizzle_MiddleOfChain) {
// struct MyStruct { a: i32; foo: vec4<f32> }
// var a: MyStruct;
// let b = a.foo.zyx.yx[0]
@@ -262,11 +262,11 @@
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, MyStruct, read_write> = var
- %1:ptr<function, vec4<f32>, read_write> = access %a, 1
- %2:vec4<f32> = load %1
- %3:vec3<f32> = swizzle %2, zxy
- %4:vec2<f32> = swizzle %3, yx
- %b:f32 = access %4, 0
+ %3:ptr<function, vec4<f32>, read_write> = access %a 1u
+ %4:vec4<f32> = load %3
+ %5:vec3<f32> = swizzle %4, zyx
+ %6:vec2<f32> = swizzle %5, yx
+ %b:f32 = access %6 0u
ret
}
}
@@ -424,7 +424,7 @@
)");
}
-TEST_F(IR_FromProgramAccessorTest, DISABLED_Accessor_Let_MultiElementSwizzle) {
+TEST_F(IR_FromProgramAccessorTest, Accessor_Let_MultiElementSwizzle) {
// let a: vec3<f32 = vec3()>
// let b = a.zyxz
@@ -438,15 +438,14 @@
EXPECT_EQ(Disassemble(m.Get()),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- %1:vec3<f32> = construct vec3<f32>
- %b:vec4<f32> = swizzle vec3<f32>(0u), zyxz
+ %b:vec4<f32> = swizzle vec3<f32>(0.0f), zyxz
ret
}
}
)");
}
-TEST_F(IR_FromProgramAccessorTest, DISABLED_Accessor_Let_MultiElementSwizzleOfSwizzle) {
+TEST_F(IR_FromProgramAccessorTest, Accessor_Let_MultiElementSwizzleOfSwizzle) {
// let a: vec3<f32> = vec3();
// let b = a.zyx.yy
@@ -460,7 +459,7 @@
EXPECT_EQ(Disassemble(m.Get()),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- %2:vec3<f32> = swizzle vec3<f32>(0u) zyx
+ %2:vec3<f32> = swizzle vec3<f32>(0.0f), zyx
%b:vec2<f32> = swizzle %2, yy
ret
}
@@ -468,7 +467,7 @@
)");
}
-TEST_F(IR_FromProgramAccessorTest, DISABLED_Accessor_Let_MultiElementSwizzle_MiddleOfChain) {
+TEST_F(IR_FromProgramAccessorTest, Accessor_Let_MultiElementSwizzle_MiddleOfChain) {
// struct MyStruct { a: i32; foo: vec4<f32> }
// let a: MyStruct = MyStruct();
// let b = a.foo.zyx.yx[0]
@@ -489,10 +488,10 @@
EXPECT_EQ(Disassemble(m.Get()),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- %2:vec4<f32> = access MyStruct(), 1u
- %2:vec3<f32> = swizzle %1, zxy
- %3:vec2<f32> = swizzle %2, yx
- %b:f32 = access %3, 0u
+ %2:vec4<f32> = access MyStruct(0i, vec4<f32>(0.0f)) 1u
+ %3:vec3<f32> = swizzle %2, zyx
+ %4:vec2<f32> = swizzle %3, yx
+ %b:f32 = access %4 0u
ret
}
}
diff --git a/src/tint/ir/swizzle.cc b/src/tint/ir/swizzle.cc
new file mode 100644
index 0000000..1887d4d
--- /dev/null
+++ b/src/tint/ir/swizzle.cc
@@ -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.
+
+#include "src/tint/ir/swizzle.h"
+
+#include <utility>
+
+#include "src/tint/debug.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::Swizzle);
+
+namespace tint::ir {
+
+Swizzle::Swizzle(const type::Type* ty, Value* object, utils::VectorRef<uint32_t> indices)
+ : result_type_(ty), object_(object), indices_(std::move(indices)) {
+ object_->AddUsage(this);
+}
+
+Swizzle::~Swizzle() = default;
+
+} // namespace tint::ir
diff --git a/src/tint/ir/swizzle.h b/src/tint/ir/swizzle.h
new file mode 100644
index 0000000..31c3da1
--- /dev/null
+++ b/src/tint/ir/swizzle.h
@@ -0,0 +1,50 @@
+// 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_SWIZZLE_H_
+#define SRC_TINT_IR_SWIZZLE_H_
+
+#include "src/tint/ir/instruction.h"
+#include "src/tint/utils/castable.h"
+
+namespace tint::ir {
+
+/// A swizzle instruction in the IR.
+class Swizzle : public utils::Castable<Swizzle, Instruction> {
+ public:
+ /// Constructor
+ /// @param result_type the result type
+ /// @param object the object being swizzled
+ /// @param indices the indices to swizzle
+ Swizzle(const type::Type* result_type, Value* object, utils::VectorRef<uint32_t> indices);
+ ~Swizzle() override;
+
+ /// @returns the type of the value
+ const type::Type* Type() const override { return result_type_; }
+
+ /// @returns the object used for the access
+ Value* Object() const { return object_; }
+
+ /// @returns the swizzle indices
+ utils::VectorRef<uint32_t> Indices() const { return indices_; }
+
+ private:
+ const type::Type* result_type_ = nullptr;
+ Value* object_ = nullptr;
+ utils::Vector<uint32_t, 4> indices_;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_SWIZZLE_H_
