encounter/dawn-cmake
1
0
Fork 0
mirror of https://github.com/encounter/dawn-cmake.git synced 2025-12-21 18:59:21 +00:00
Code Issues Packages Projects Releases Wiki Activity

Tint/transform: make AddBlockAttribute always do wrapping if possible

Browse Source 
This CL make transform AddBlockAttribute always try to wrap types used
by buffer variables into a struct, in order to generate valid GLSL code
for assigning one buffer struct variable to another buffer struct
variable.

Fixed: tint:1735
Change-Id: I009d8a9ca7ecea1dc0ad6164275c964a18acb33f
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/108023
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Zhaoming Jiang <zhaoming.jiang@intel.com>
This commit is contained in:
Zhaoming Jiang
2022-11-02 02:25:38 +00:00
committed by Dawn LUCI CQ
parent 2bea9055f4
commit 6ab5d3c151
465 changed files with 15681 additions and 12853 deletions
Download Patch File Download Diff File Expand all files Collapse all files

46
src/tint/transform/add_block_attribute.cc
View File

@@ -27,19 +27,6 @@ TINT_INSTANTIATE_TYPEINFO(tint::transform::AddBlockAttribute::BlockAttribute);
namespace tint::transform {
namespace {
bool IsUsedAsNonBuffer(const std::unordered_set<tint::ast::AddressSpace>& uses) {
for (auto use : uses) {
if (!ast::IsHostShareable(use)) {
return true;
}
}
return false;
}
} // namespace
AddBlockAttribute::AddBlockAttribute() = default;
AddBlockAttribute::~AddBlockAttribute() = default;
@@ -47,25 +34,6 @@ AddBlockAttribute::~AddBlockAttribute() = default;
void AddBlockAttribute::Run(CloneContext& ctx, const DataMap&, DataMap&) const {
auto& sem = ctx.src->Sem();
// Collect the set of structs that are nested in other types.
utils::Hashset<const sem::Struct*, 8> nested_structs;
for (auto* ty : ctx.src->Types()) {
Switch(
ty,
[&](const sem::Array* arr) {
if (auto* nested_str = arr->ElemType()->As<sem::Struct>()) {
nested_structs.Add(nested_str);
}
},
[&](const sem::Struct* str) {
for (auto* member : str->Members()) {
if (auto* nested_str = member->Type()->As<sem::Struct>()) {
nested_structs.Add(nested_str);
}
}
});
}
// A map from a type in the source program to a block-decorated wrapper that contains it in the
// destination program.
utils::Hashmap<const sem::Type*, const ast::Struct*, 8> wrapper_structs;
@@ -80,16 +48,18 @@ void AddBlockAttribute::Run(CloneContext& ctx, const DataMap&, DataMap&) const {
auto* ty = var->Type()->UnwrapRef();
auto* str = ty->As<sem::Struct>();
bool needs_wrapping =
!str || // Type is not a structure
nested_structs.Contains(str) || // Structure is nested by another type
IsUsedAsNonBuffer(str->AddressSpaceUsage()); // Structure is used as a non-buffer usage
// Always try to wrap the buffer type into a struct. We can not do so only if it is a struct
// but without a fixed footprint, i.e. contains a runtime-sized array as its member. Note
// that such struct type can be only used as storage buffer variables' type. Also note that
// any buffer struct type that may be nested by another type must have a fixed footprint,
// therefore will be wrapped.
bool needs_wrapping = !str || // Type is not a structure
str->HasFixedFootprint(); // Struct has a fixed footprint
if (needs_wrapping) {
const char* kMemberName = "inner";
// This is a non-struct or a struct that is nested somewhere else, so we
// need to wrap it first.
auto* wrapper = wrapper_structs.GetOrCreate(ty, [&] {
auto* block = ctx.dst->ASTNodes().Create<BlockAttribute>(ctx.dst->ID(),
ctx.dst->AllocateNodeID());

7
src/tint/transform/add_block_attribute.h
View File

@@ -22,11 +22,8 @@
namespace tint::transform {
/// AddBlockAttribute is a transform that adds an
/// `@internal(block)` attribute to any structure that is used as the
/// store type of a buffer. If that structure is nested inside another structure
/// or an array, then it is wrapped inside another structure which gets the
/// `@internal(block)` attribute instead.
/// AddBlockAttribute is a transform that wrap the store type of a buffer into a struct if possible,
/// then adds an `@internal(block)` attribute to the wrapper struct.
class AddBlockAttribute final : public Castable<AddBlockAttribute, Transform> {
public:
/// BlockAttribute is an InternalAttribute that is used to decorate a

221
src/tint/transform/add_block_attribute_test.cc
View File

@@ -200,6 +200,85 @@ fn main() {
EXPECT_EQ(expect, str(got));
}
TEST_F(AddBlockAttributeTest, BasicStruct_Storage_AccessRoot) {
auto* src = R"(
struct S {
f : f32,
};
@group(0) @binding(0)
var<storage, read_write> s : S;
@fragment
fn main() {
let f = s;
}
)";
auto* expect = R"(
struct S {
f : f32,
}
@internal(block)
struct s_block {
inner : S,
}
@group(0) @binding(0) var<storage, read_write> s : s_block;
@fragment
fn main() {
let f = s.inner;
}
)";
auto got = Run<AddBlockAttribute>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(AddBlockAttributeTest, BasicStruct_Storage_TwoUsage_AccessRoot) {
auto* src = R"(
struct S {
f : f32,
};
@group(0) @binding(0)
var<storage, read_write> in : S;
@group(0) @binding(1)
var<storage, read_write> out : S;
@compute @workgroup_size(1)
fn main() {
out = in;
}
)";
auto* expect = R"(
struct S {
f : f32,
}
@internal(block)
struct in_block {
inner : S,
}
@group(0) @binding(0) var<storage, read_write> in : in_block;
@group(0) @binding(1) var<storage, read_write> out : in_block;
@compute @workgroup_size(1)
fn main() {
out.inner = in.inner;
}
)";
auto got = Run<AddBlockAttribute>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(AddBlockAttributeTest, BasicStruct_AccessField) {
auto* src = R"(
struct S {
@@ -215,16 +294,20 @@ fn main() {
}
)";
auto* expect = R"(
@internal(block)
struct S {
f : f32,
}
@group(0) @binding(0) var<uniform> u : S;
@internal(block)
struct u_block {
inner : S,
}
@group(0) @binding(0) var<uniform> u : u_block;
@fragment
fn main() {
let f = u.f;
let f = u.inner.f;
}
)";
@@ -280,16 +363,20 @@ fn main() {
auto* expect = R"(
enable chromium_experimental_push_constant;
@internal(block)
struct S {
f : f32,
}
var<push_constant> u : S;
@internal(block)
struct u_block {
inner : S,
}
var<push_constant> u : u_block;
@fragment
fn main() {
let f = u.f;
let f = u.inner.f;
}
)";
@@ -321,16 +408,20 @@ struct Inner {
f : f32,
}
@internal(block)
struct Outer {
i : Inner,
}
@group(0) @binding(0) var<uniform> u : Outer;
@internal(block)
struct u_block {
inner : Outer,
}
@group(0) @binding(0) var<uniform> u : u_block;
@fragment
fn main() {
let f = u.i.f;
let f = u.inner.i.f;
}
)";
@@ -366,12 +457,16 @@ struct Inner {
f : f32,
}
@internal(block)
struct Outer {
i : Inner,
}
@group(0) @binding(0) var<uniform> u0 : Outer;
@internal(block)
struct u0_block {
inner : Outer,
}
@group(0) @binding(0) var<uniform> u0 : u0_block;
@internal(block)
struct u1_block {
@@ -382,7 +477,7 @@ struct u1_block {
@fragment
fn main() {
let f0 = u0.i.f;
let f0 = u0.inner.i.f;
let f1 = u1.inner.f;
}
)";
@@ -474,12 +569,16 @@ struct Inner {
f : f32,
}
@internal(block)
struct S {
i : Inner,
}
@group(0) @binding(0) var<uniform> u0 : S;
@internal(block)
struct u0_block {
inner : S,
}
@group(0) @binding(0) var<uniform> u0 : u0_block;
@internal(block)
struct u1_block {
@@ -492,7 +591,7 @@ struct u1_block {
@fragment
fn main() {
let f0 = u0.i.f;
let f0 = u0.inner.i.f;
let f1 = u1.inner.f;
let f2 = u2.inner.f;
}
@@ -621,14 +720,18 @@ struct Inner {
type MyInner = Inner;
@internal(block)
struct Outer {
i : MyInner,
}
type MyOuter = Outer;
@group(0) @binding(0) var<uniform> u0 : MyOuter;
@internal(block)
struct u0_block {
inner : Outer,
}
@group(0) @binding(0) var<uniform> u0 : u0_block;
@internal(block)
struct u1_block {
@@ -639,7 +742,7 @@ struct u1_block {
@fragment
fn main() {
let f0 = u0.i.f;
let f0 = u0.inner.i.f;
let f1 = u1.inner.f;
}
)";
@@ -678,7 +781,7 @@ struct Inner {
auto* expect = R"(
@fragment
fn main() {
let f0 = u0.i.f;
let f0 = u0.inner.i.f;
let f1 = u1.inner.f;
}
@@ -691,11 +794,15 @@ struct u1_block {
type MyInner = Inner;
@group(0) @binding(0) var<uniform> u0 : MyOuter;
@internal(block)
struct u0_block {
inner : Outer,
}
@group(0) @binding(0) var<uniform> u0 : u0_block;
type MyOuter = Outer;
@internal(block)
struct Outer {
i : MyInner,
}
@@ -810,18 +917,22 @@ fn main() {
}
)";
auto* expect = R"(
@internal(block) @internal(block)
struct S {
f : f32,
}
@group(0) @binding(0) var<uniform> u : S;
@internal(block)
struct u_block {
inner : S,
}
@group(0) @binding(1) var<storage, read_write> s : S;
@group(0) @binding(0) var<uniform> u : u_block;
@group(0) @binding(1) var<storage, read_write> s : u_block;
@fragment
fn main() {
s = u;
s.inner = u.inner;
}
)";
@@ -850,5 +961,65 @@ fn main() {
EXPECT_EQ(expect, str(got));
}
TEST_F(AddBlockAttributeTest, StorageBufferWithRuntimeArray) {
auto* src = R"(
struct S {
f : f32,
}
struct SWithArr {
f : f32,
arr : array<f32>,
}
@group(0) @binding(0)
var<storage, read> in_1 : S;
@group(0) @binding(1)
var<storage, read> in_2 : SWithArr;
@group(1) @binding(0)
var<storage, read_write> out : SWithArr;
@fragment
fn main() {
out.f = in_1.f;
out.arr[0] = in_2.arr[1];
}
)";
auto* expect = R"(
struct S {
f : f32,
}
@internal(block) @internal(block)
struct SWithArr {
f : f32,
arr : array<f32>,
}
@internal(block)
struct in_1_block {
inner : S,
}
@group(0) @binding(0) var<storage, read> in_1 : in_1_block;
@group(0) @binding(1) var<storage, read> in_2 : SWithArr;
@group(1) @binding(0) var<storage, read_write> out : SWithArr;
@fragment
fn main() {
out.f = in_1.inner.f;
out.arr[0] = in_2.arr[1];
}
)";
auto got = Run<AddBlockAttribute>(src);
EXPECT_EQ(expect, str(got));
}
} // namespace
} // namespace tint::transform