inspector: Add support for non-struct buffers
Use the sem::Type size instead of the sem::Struct size.
Some binding numbers in MSL expected outputs changed since these
buffers were previously not being reported in the resource list that
the Tint exe uses to automatically remap bindings.
Bug: tint:1372
Change-Id: I14479ddc5129d91b91449cc2d68ee37bd99d2f7e
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/76220
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/inspector/inspector.cc b/src/inspector/inspector.cc
index 33c8449..1b4cede 100644
--- a/src/inspector/inspector.cc
+++ b/src/inspector/inspector.cc
@@ -317,16 +317,10 @@
size_t size = 0;
auto* func_sem = program_->Sem().Get(func);
for (auto& ruv : func_sem->TransitivelyReferencedUniformVariables()) {
- const sem::Struct* s = ruv.first->Type()->UnwrapRef()->As<sem::Struct>();
- if (s) {
- size += s->Size();
- }
+ size += ruv.first->Type()->UnwrapRef()->Size();
}
for (auto& rsv : func_sem->TransitivelyReferencedStorageBufferVariables()) {
- const sem::Struct* s = rsv.first->Type()->UnwrapRef()->As<sem::Struct>();
- if (s) {
- size += s->Size();
- }
+ size += rsv.first->Type()->UnwrapRef()->Size();
}
if (static_cast<uint64_t>(size) >
@@ -377,17 +371,18 @@
auto binding_info = ruv.second;
auto* unwrapped_type = var->Type()->UnwrapRef();
- auto* str = unwrapped_type->As<sem::Struct>();
- if (str == nullptr) {
- continue;
- }
ResourceBinding entry;
entry.resource_type = ResourceBinding::ResourceType::kUniformBuffer;
entry.bind_group = binding_info.group->value;
entry.binding = binding_info.binding->value;
- entry.size = str->Size();
- entry.size_no_padding = str->SizeNoPadding();
+ entry.size = unwrapped_type->Size();
+ entry.size_no_padding = entry.size;
+ if (auto* str = unwrapped_type->As<sem::Struct>()) {
+ entry.size_no_padding = str->SizeNoPadding();
+ } else {
+ entry.size_no_padding = entry.size;
+ }
result.push_back(entry);
}
@@ -667,10 +662,7 @@
continue;
}
- auto* str = var->Type()->UnwrapRef()->As<sem::Struct>();
- if (!str) {
- continue;
- }
+ auto* unwrapped_type = var->Type()->UnwrapRef();
ResourceBinding entry;
entry.resource_type =
@@ -678,8 +670,12 @@
: ResourceBinding::ResourceType::kStorageBuffer;
entry.bind_group = binding_info.group->value;
entry.binding = binding_info.binding->value;
- entry.size = str->Size();
- entry.size_no_padding = str->SizeNoPadding();
+ entry.size = unwrapped_type->Size();
+ if (auto* str = unwrapped_type->As<sem::Struct>()) {
+ entry.size_no_padding = str->SizeNoPadding();
+ } else {
+ entry.size_no_padding = entry.size;
+ }
result.push_back(entry);
}
diff --git a/src/inspector/inspector_test.cc b/src/inspector/inspector_test.cc
index 1c2d6da..1d237f8 100644
--- a/src/inspector/inspector_test.cc
+++ b/src/inspector/inspector_test.cc
@@ -1218,7 +1218,24 @@
EXPECT_EQ(0u, inspector.GetStorageSize("ep_func"));
}
-TEST_F(InspectorGetStorageSizeTest, Simple) {
+TEST_F(InspectorGetStorageSizeTest, Simple_NonStruct) {
+ AddUniformBuffer("ub_var", ty.i32(), 0, 0);
+ AddStorageBuffer("sb_var", ty.i32(), ast::Access::kReadWrite, 1, 0);
+ AddStorageBuffer("rosb_var", ty.i32(), ast::Access::kRead, 1, 1);
+ Func("ep_func", {}, ty.void_(),
+ {
+ Decl(Const("ub", nullptr, Expr("ub_var"))),
+ Decl(Const("sb", nullptr, Expr("sb_var"))),
+ Decl(Const("rosb", nullptr, Expr("rosb_var"))),
+ },
+ {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
+
+ Inspector& inspector = Build();
+
+ EXPECT_EQ(12u, inspector.GetStorageSize("ep_func"));
+}
+
+TEST_F(InspectorGetStorageSizeTest, Simple_Struct) {
auto* ub_struct_type = MakeUniformBufferType("ub_type", {ty.i32(), ty.i32()});
AddUniformBuffer("ub_var", ty.Of(ub_struct_type), 0, 0);
MakeStructVariableReferenceBodyFunction("ub_func", "ub_var", {{0, ty.i32()}});
@@ -1243,6 +1260,33 @@
EXPECT_EQ(16u, inspector.GetStorageSize("ep_func"));
}
+TEST_F(InspectorGetStorageSizeTest, NonStructVec3) {
+ AddUniformBuffer("ub_var", ty.vec3<f32>(), 0, 0);
+ Func("ep_func", {}, ty.void_(),
+ {
+ Decl(Const("ub", nullptr, Expr("ub_var"))),
+ },
+ {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
+
+ Inspector& inspector = Build();
+
+ EXPECT_EQ(12u, inspector.GetStorageSize("ep_func"));
+}
+
+TEST_F(InspectorGetStorageSizeTest, StructVec3) {
+ auto* ub_struct_type = MakeUniformBufferType("ub_type", {ty.vec3<f32>()});
+ AddUniformBuffer("ub_var", ty.Of(ub_struct_type), 0, 0);
+ Func("ep_func", {}, ty.void_(),
+ {
+ Decl(Const("ub", nullptr, Expr("ub_var"))),
+ },
+ {Stage(ast::PipelineStage::kCompute), WorkgroupSize(1)});
+
+ Inspector& inspector = Build();
+
+ EXPECT_EQ(16u, inspector.GetStorageSize("ep_func"));
+}
+
TEST_F(InspectorGetResourceBindingsTest, Empty) {
MakeCallerBodyFunction("ep_func", {},
ast::DecorationList{
@@ -1381,7 +1425,30 @@
EXPECT_TRUE(error.find("not an entry point") != std::string::npos);
}
-TEST_F(InspectorGetUniformBufferResourceBindingsTest, Simple) {
+TEST_F(InspectorGetUniformBufferResourceBindingsTest, Simple_NonStruct) {
+ AddUniformBuffer("foo_ub", ty.i32(), 0, 0);
+ MakePlainGlobalReferenceBodyFunction("ub_func", "foo_ub", ty.i32(), {});
+
+ MakeCallerBodyFunction("ep_func", {"ub_func"},
+ ast::DecorationList{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ Inspector& inspector = Build();
+
+ auto result = inspector.GetUniformBufferResourceBindings("ep_func");
+ ASSERT_FALSE(inspector.has_error()) << inspector.error();
+ ASSERT_EQ(1u, result.size());
+
+ EXPECT_EQ(ResourceBinding::ResourceType::kUniformBuffer,
+ result[0].resource_type);
+ EXPECT_EQ(0u, result[0].bind_group);
+ EXPECT_EQ(0u, result[0].binding);
+ EXPECT_EQ(4u, result[0].size);
+ EXPECT_EQ(4u, result[0].size_no_padding);
+}
+
+TEST_F(InspectorGetUniformBufferResourceBindingsTest, Simple_Struct) {
auto* foo_struct_type = MakeUniformBufferType("foo_type", {ty.i32()});
AddUniformBuffer("foo_ub", ty.Of(foo_struct_type), 0, 0);
@@ -1459,6 +1526,29 @@
EXPECT_EQ(12u, result[0].size_no_padding);
}
+TEST_F(InspectorGetUniformBufferResourceBindingsTest, NonStructVec3) {
+ AddUniformBuffer("foo_ub", ty.vec3<f32>(), 0, 0);
+ MakePlainGlobalReferenceBodyFunction("ub_func", "foo_ub", ty.vec3<f32>(), {});
+
+ MakeCallerBodyFunction("ep_func", {"ub_func"},
+ ast::DecorationList{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ Inspector& inspector = Build();
+
+ auto result = inspector.GetUniformBufferResourceBindings("ep_func");
+ ASSERT_FALSE(inspector.has_error()) << inspector.error();
+ ASSERT_EQ(1u, result.size());
+
+ EXPECT_EQ(ResourceBinding::ResourceType::kUniformBuffer,
+ result[0].resource_type);
+ EXPECT_EQ(0u, result[0].bind_group);
+ EXPECT_EQ(0u, result[0].binding);
+ EXPECT_EQ(12u, result[0].size);
+ EXPECT_EQ(12u, result[0].size_no_padding);
+}
+
TEST_F(InspectorGetUniformBufferResourceBindingsTest, MultipleUniformBuffers) {
auto* ub_struct_type =
MakeUniformBufferType("ub_type", {ty.i32(), ty.u32(), ty.f32()});
@@ -1546,7 +1636,30 @@
EXPECT_EQ(80u, result[0].size_no_padding);
}
-TEST_F(InspectorGetStorageBufferResourceBindingsTest, Simple) {
+TEST_F(InspectorGetStorageBufferResourceBindingsTest, Simple_NonStruct) {
+ AddStorageBuffer("foo_sb", ty.i32(), ast::Access::kReadWrite, 0, 0);
+ MakePlainGlobalReferenceBodyFunction("sb_func", "foo_sb", ty.i32(), {});
+
+ MakeCallerBodyFunction("ep_func", {"sb_func"},
+ ast::DecorationList{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ Inspector& inspector = Build();
+
+ auto result = inspector.GetStorageBufferResourceBindings("ep_func");
+ ASSERT_FALSE(inspector.has_error()) << inspector.error();
+ ASSERT_EQ(1u, result.size());
+
+ EXPECT_EQ(ResourceBinding::ResourceType::kStorageBuffer,
+ result[0].resource_type);
+ EXPECT_EQ(0u, result[0].bind_group);
+ EXPECT_EQ(0u, result[0].binding);
+ EXPECT_EQ(4u, result[0].size);
+ EXPECT_EQ(4u, result[0].size_no_padding);
+}
+
+TEST_F(InspectorGetStorageBufferResourceBindingsTest, Simple_Struct) {
auto foo_struct_type = MakeStorageBufferTypes("foo_type", {ty.i32()});
AddStorageBuffer("foo_sb", foo_struct_type(), ast::Access::kReadWrite, 0, 0);
@@ -1743,6 +1856,29 @@
EXPECT_EQ(12u, result[0].size_no_padding);
}
+TEST_F(InspectorGetStorageBufferResourceBindingsTest, NonStructVec3) {
+ AddStorageBuffer("foo_ub", ty.vec3<f32>(), ast::Access::kReadWrite, 0, 0);
+ MakePlainGlobalReferenceBodyFunction("ub_func", "foo_ub", ty.vec3<f32>(), {});
+
+ MakeCallerBodyFunction("ep_func", {"ub_func"},
+ ast::DecorationList{
+ Stage(ast::PipelineStage::kFragment),
+ });
+
+ Inspector& inspector = Build();
+
+ auto result = inspector.GetStorageBufferResourceBindings("ep_func");
+ ASSERT_FALSE(inspector.has_error()) << inspector.error();
+ ASSERT_EQ(1u, result.size());
+
+ EXPECT_EQ(ResourceBinding::ResourceType::kStorageBuffer,
+ result[0].resource_type);
+ EXPECT_EQ(0u, result[0].bind_group);
+ EXPECT_EQ(0u, result[0].binding);
+ EXPECT_EQ(12u, result[0].size);
+ EXPECT_EQ(12u, result[0].size_no_padding);
+}
+
TEST_F(InspectorGetStorageBufferResourceBindingsTest, SkipReadOnly) {
auto foo_struct_type = MakeStorageBufferTypes("foo_type", {ty.i32()});
AddStorageBuffer("foo_sb", foo_struct_type(), ast::Access::kRead, 0, 0);
diff --git a/test/buffer/storage/types/array.wgsl.expected.msl b/test/buffer/storage/types/array.wgsl.expected.msl
index a95963f..685b94a 100644
--- a/test/buffer/storage/types/array.wgsl.expected.msl
+++ b/test/buffer/storage/types/array.wgsl.expected.msl
@@ -5,7 +5,7 @@
/* 0x0000 */ float arr[4];
};
-kernel void tint_symbol(device tint_array_wrapper* tint_symbol_1 [[buffer(1)]], const device tint_array_wrapper* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device tint_array_wrapper* tint_symbol_1 [[buffer(0)]], const device tint_array_wrapper* tint_symbol_2 [[buffer(1)]]) {
*(tint_symbol_1) = *(tint_symbol_2);
return;
}
diff --git a/test/buffer/storage/types/f32.wgsl.expected.msl b/test/buffer/storage/types/f32.wgsl.expected.msl
index 033b5e2..c4b1753 100644
--- a/test/buffer/storage/types/f32.wgsl.expected.msl
+++ b/test/buffer/storage/types/f32.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol(device float* tint_symbol_1 [[buffer(1)]], const device float* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device float* tint_symbol_1 [[buffer(0)]], const device float* tint_symbol_2 [[buffer(1)]]) {
*(tint_symbol_1) = *(tint_symbol_2);
return;
}
diff --git a/test/buffer/storage/types/i32.wgsl.expected.msl b/test/buffer/storage/types/i32.wgsl.expected.msl
index 1de9e03..d801b78 100644
--- a/test/buffer/storage/types/i32.wgsl.expected.msl
+++ b/test/buffer/storage/types/i32.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol(device int* tint_symbol_1 [[buffer(1)]], const device int* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device int* tint_symbol_1 [[buffer(0)]], const device int* tint_symbol_2 [[buffer(1)]]) {
*(tint_symbol_1) = *(tint_symbol_2);
return;
}
diff --git a/test/buffer/storage/types/mat2x2.wgsl.expected.msl b/test/buffer/storage/types/mat2x2.wgsl.expected.msl
index c9ffff7..fbc66bf 100644
--- a/test/buffer/storage/types/mat2x2.wgsl.expected.msl
+++ b/test/buffer/storage/types/mat2x2.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol(device float2x2* tint_symbol_1 [[buffer(1)]], const device float2x2* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device float2x2* tint_symbol_1 [[buffer(0)]], const device float2x2* tint_symbol_2 [[buffer(1)]]) {
*(tint_symbol_1) = *(tint_symbol_2);
return;
}
diff --git a/test/buffer/storage/types/mat2x3.wgsl.expected.msl b/test/buffer/storage/types/mat2x3.wgsl.expected.msl
index 1b704d4..16b1298 100644
--- a/test/buffer/storage/types/mat2x3.wgsl.expected.msl
+++ b/test/buffer/storage/types/mat2x3.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol(device float2x3* tint_symbol_1 [[buffer(1)]], const device float2x3* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device float2x3* tint_symbol_1 [[buffer(0)]], const device float2x3* tint_symbol_2 [[buffer(1)]]) {
*(tint_symbol_1) = *(tint_symbol_2);
return;
}
diff --git a/test/buffer/storage/types/mat3x2.wgsl.expected.msl b/test/buffer/storage/types/mat3x2.wgsl.expected.msl
index b8765f0..b74b79a 100644
--- a/test/buffer/storage/types/mat3x2.wgsl.expected.msl
+++ b/test/buffer/storage/types/mat3x2.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol(device float3x2* tint_symbol_1 [[buffer(1)]], const device float3x2* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device float3x2* tint_symbol_1 [[buffer(0)]], const device float3x2* tint_symbol_2 [[buffer(1)]]) {
*(tint_symbol_1) = *(tint_symbol_2);
return;
}
diff --git a/test/buffer/storage/types/mat4x4.wgsl.expected.msl b/test/buffer/storage/types/mat4x4.wgsl.expected.msl
index 6b33874..d52f1c4 100644
--- a/test/buffer/storage/types/mat4x4.wgsl.expected.msl
+++ b/test/buffer/storage/types/mat4x4.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol(device float4x4* tint_symbol_1 [[buffer(1)]], const device float4x4* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device float4x4* tint_symbol_1 [[buffer(0)]], const device float4x4* tint_symbol_2 [[buffer(1)]]) {
*(tint_symbol_1) = *(tint_symbol_2);
return;
}
diff --git a/test/buffer/storage/types/runtime_array.wgsl.expected.msl b/test/buffer/storage/types/runtime_array.wgsl.expected.msl
index ff6a283..573d6ce 100644
--- a/test/buffer/storage/types/runtime_array.wgsl.expected.msl
+++ b/test/buffer/storage/types/runtime_array.wgsl.expected.msl
@@ -11,7 +11,7 @@
/* 0x0000 */ S arr[1];
};
-kernel void tint_symbol(device tint_symbol_2* tint_symbol_1 [[buffer(1)]], const device tint_symbol_4* tint_symbol_3 [[buffer(0)]]) {
+kernel void tint_symbol(device tint_symbol_2* tint_symbol_1 [[buffer(0)]], const device tint_symbol_4* tint_symbol_3 [[buffer(1)]]) {
(*(tint_symbol_1)).arr[0] = (*(tint_symbol_3)).arr[0];
return;
}
diff --git a/test/buffer/storage/types/u32.wgsl.expected.msl b/test/buffer/storage/types/u32.wgsl.expected.msl
index bb65310..c2e500a 100644
--- a/test/buffer/storage/types/u32.wgsl.expected.msl
+++ b/test/buffer/storage/types/u32.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol(device uint* tint_symbol_1 [[buffer(1)]], const device uint* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device uint* tint_symbol_1 [[buffer(0)]], const device uint* tint_symbol_2 [[buffer(1)]]) {
*(tint_symbol_1) = *(tint_symbol_2);
return;
}
diff --git a/test/buffer/storage/types/vec2.wgsl.expected.msl b/test/buffer/storage/types/vec2.wgsl.expected.msl
index c730e05..6627fa2 100644
--- a/test/buffer/storage/types/vec2.wgsl.expected.msl
+++ b/test/buffer/storage/types/vec2.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol(device int2* tint_symbol_1 [[buffer(1)]], const device int2* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device int2* tint_symbol_1 [[buffer(0)]], const device int2* tint_symbol_2 [[buffer(1)]]) {
*(tint_symbol_1) = *(tint_symbol_2);
return;
}
diff --git a/test/buffer/storage/types/vec3.wgsl.expected.msl b/test/buffer/storage/types/vec3.wgsl.expected.msl
index 6d55796..54cc7c8 100644
--- a/test/buffer/storage/types/vec3.wgsl.expected.msl
+++ b/test/buffer/storage/types/vec3.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol(device uint3* tint_symbol_1 [[buffer(1)]], const device uint3* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device uint3* tint_symbol_1 [[buffer(0)]], const device uint3* tint_symbol_2 [[buffer(1)]]) {
*(tint_symbol_1) = *(tint_symbol_2);
return;
}
diff --git a/test/buffer/storage/types/vec4.wgsl.expected.msl b/test/buffer/storage/types/vec4.wgsl.expected.msl
index 4af349b..65cbadb 100644
--- a/test/buffer/storage/types/vec4.wgsl.expected.msl
+++ b/test/buffer/storage/types/vec4.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-kernel void tint_symbol(device float4* tint_symbol_1 [[buffer(1)]], const device float4* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(device float4* tint_symbol_1 [[buffer(0)]], const device float4* tint_symbol_2 [[buffer(1)]]) {
*(tint_symbol_1) = *(tint_symbol_2);
return;
}
