[ir] Emit struct declarations in the disassembly
Bug: tint:1718
Change-Id: Iff88d362c20381280ec98f868107a1ecd144c01f
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/136260
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 1d6b421..31e3868 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -42,6 +42,7 @@
#include "src/tint/ir/user_call.h"
#include "src/tint/ir/var.h"
#include "src/tint/switch.h"
+#include "src/tint/type/struct.h"
#include "src/tint/type/type.h"
#include "src/tint/utils/scoped_assignment.h"
#include "src/tint/utils/string.h"
@@ -95,6 +96,12 @@
}
std::string Disassembler::Disassemble() {
+ for (auto* ty : mod_.Types()) {
+ if (auto* str = ty->As<type::Struct>()) {
+ EmitStructDecl(str);
+ }
+ }
+
if (mod_.root_block) {
Indent() << "# Root block" << std::endl;
WalkInternal(mod_.root_block);
@@ -669,4 +676,33 @@
out_ << std::endl;
}
+void Disassembler::EmitStructDecl(const type::Struct* str) {
+ out_ << str->Name().Name() << " = struct @align(" << str->Align() << ") {" << std::endl;
+ for (auto* member : str->Members()) {
+ out_ << " " << member->Name().Name() << ":" << member->Type()->FriendlyName();
+ out_ << " @offset(" << member->Offset() << ")";
+ if (member->Attributes().invariant) {
+ out_ << ", @invariant";
+ }
+ if (member->Attributes().location.has_value()) {
+ out_ << ", @location(" << member->Attributes().location.value() << ")";
+ }
+ if (member->Attributes().interpolation.has_value()) {
+ auto& interp = member->Attributes().interpolation.value();
+ out_ << ", @interpolate(" << interp.type;
+ if (interp.sampling != builtin::InterpolationSampling::kUndefined) {
+ out_ << ", " << interp.sampling;
+ }
+ out_ << ")";
+ }
+ if (member->Attributes().builtin.has_value()) {
+ out_ << ", @builtin(" << member->Attributes().builtin.value() << ")";
+ }
+
+ out_ << std::endl;
+ }
+ out_ << "}" << std::endl;
+ out_ << std::endl;
+}
+
} // namespace tint::ir
diff --git a/src/tint/ir/disassembler.h b/src/tint/ir/disassembler.h
index bf2f8fa..ea27470 100644
--- a/src/tint/ir/disassembler.h
+++ b/src/tint/ir/disassembler.h
@@ -29,6 +29,11 @@
#include "src/tint/utils/hashset.h"
#include "src/tint/utils/string_stream.h"
+// Forward declarations.
+namespace tint::type {
+class Struct;
+}
+
namespace tint::ir {
/// Helper class to disassemble the IR
@@ -74,6 +79,7 @@
void EmitSwitch(const Switch* s);
void EmitLoop(const Loop* l);
void EmitIf(const If* i);
+ void EmitStructDecl(const type::Struct* str);
const Module& mod_;
utils::StringStream out_;
diff --git a/src/tint/ir/from_program_accessor_test.cc b/src/tint/ir/from_program_accessor_test.cc
index 9da49a1..d8e6601 100644
--- a/src/tint/ir/from_program_accessor_test.cc
+++ b/src/tint/ir/from_program_accessor_test.cc
@@ -90,7 +90,11 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ R"(MyStruct = struct @align(4) {
+ foo:i32 @offset(0)
+}
+
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, MyStruct, read_write> = var
%3:ptr<function, i32, read_write> = access %a, 0u
@@ -122,7 +126,16 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ R"(Inner = struct @align(4) {
+ bar:f32 @offset(0)
+}
+
+Outer = struct @align(4) {
+ a:i32 @offset(0)
+ foo:Inner @offset(4)
+}
+
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, Outer, read_write> = var
%3:ptr<function, f32, read_write> = access %a, 1u, 0u
@@ -158,7 +171,18 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ R"(Inner = struct @align(16) {
+ b:i32 @offset(0)
+ c:f32 @offset(4)
+ bar:vec4<f32> @offset(16)
+}
+
+Outer = struct @align(16) {
+ a:i32 @offset(0)
+ foo:array<Inner, 4> @offset(16)
+}
+
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, array<Outer, 4>, read_write> = var
%3:ptr<function, vec4<f32>, read_write> = access %a, 0u, 1u, 1u, 2u
@@ -281,7 +305,12 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ R"(MyStruct = struct @align(16) {
+ a:i32 @offset(0)
+ foo:vec4<f32> @offset(16)
+}
+
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%a:ptr<function, MyStruct, read_write> = var
%3:ptr<function, vec4<f32>, read_write> = access %a, 1u
@@ -352,7 +381,11 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ R"(MyStruct = struct @align(4) {
+ foo:i32 @offset(0)
+}
+
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%b:i32 = access MyStruct(0i), 0u
ret
@@ -382,7 +415,16 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ R"(Inner = struct @align(4) {
+ bar:f32 @offset(0)
+}
+
+Outer = struct @align(4) {
+ a:i32 @offset(0)
+ foo:Inner @offset(4)
+}
+
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%b:f32 = access Outer(0i, Inner(0.0f)), 1u, 0u
ret
@@ -416,7 +458,18 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ R"(Inner = struct @align(16) {
+ b:i32 @offset(0)
+ c:f32 @offset(4)
+ bar:vec4<f32> @offset(16)
+}
+
+Outer = struct @align(16) {
+ a:i32 @offset(0)
+ foo:array<Inner, 4> @offset(16)
+}
+
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%b:vec4<f32> = access array<Outer, 4>(Outer(0i, array<Inner, 4>(Inner(0i, 0.0f, vec4<f32>(0.0f))))), 0u, 1u, 1u, 2u
ret
@@ -508,7 +561,12 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
- R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ R"(MyStruct = struct @align(16) {
+ a:i32 @offset(0)
+ foo:vec4<f32> @offset(16)
+}
+
+%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
%2:vec4<f32> = access MyStruct(0i, vec4<f32>(0.0f)), 1u
%3:vec3<f32> = swizzle %2, zyx
diff --git a/src/tint/ir/module.h b/src/tint/ir/module.h
index 1520888..caffbfc 100644
--- a/src/tint/ir/module.h
+++ b/src/tint/ir/module.h
@@ -67,6 +67,9 @@
Symbol SetName(const Value* value, std::string_view name);
/// @return the type manager for the module
+ const type::Manager& Types() const { return constant_values.types; }
+
+ /// @return the type manager for the module
type::Manager& Types() { return constant_values.types; }
/// The block allocator
