spirv-reader: preserve NonWritable decoration on module-scope vars
Fixed: 348129253
Change-Id: Ic48e6c8e3b39ae82718ee59c2b024625deabc2f0
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/194620
Commit-Queue: David Neto <dneto@google.com>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: David Neto <dneto@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc b/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
index efb14a2..8e486f5 100644
--- a/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/ast_parser.cc
@@ -1534,14 +1534,14 @@
// here.)
ast_initializer = MakeConstantExpression(var.GetSingleWordInOperand(1)).expr;
}
- auto ast_access = VarAccess(ast_store_type, ast_address_space);
- auto* ast_var = MakeVar(var.result_id(), ast_address_space, ast_access, ast_store_type,
- ast_initializer, Attributes{});
+ auto* ast_var = MakeVar(var.result_id(), ast_address_space, ast_store_type, ast_initializer,
+ Attributes{});
// TODO(dneto): initializers (a.k.a. initializer expression)
if (ast_var) {
builder_.AST().AddGlobalVariable(ast_var);
module_variable_.GetOrAdd(var.result_id(), [&] {
- return ModuleVariable{ast_var, ast_address_space, ast_access};
+ auto access = VarAccess(var.result_id(), ast_store_type, ast_address_space);
+ return ModuleVariable{ast_var, ast_address_space, access};
});
}
}
@@ -1571,8 +1571,7 @@
}
auto storage_type = ConvertType(builtin_position_.position_member_type_id);
auto ast_address_space = enum_converter_.ToAddressSpace(builtin_position_.storage_class);
- auto ast_access = VarAccess(storage_type, ast_address_space);
- auto* ast_var = MakeVar(builtin_position_.per_vertex_var_id, ast_address_space, ast_access,
+ auto* ast_var = MakeVar(builtin_position_.per_vertex_var_id, ast_address_space,
storage_type, ast_initializer, {});
builder_.AST().AddGlobalVariable(ast_var);
@@ -1606,14 +1605,16 @@
return size->AsIntConstant();
}
-core::Access ASTParser::VarAccess(const Type* storage_type, core::AddressSpace address_space) {
+core::Access ASTParser::VarAccess(uint32_t var_id,
+ const Type* storage_type,
+ core::AddressSpace address_space) {
if (address_space != core::AddressSpace::kStorage) {
return core::Access::kUndefined;
}
- bool read_only = false;
+ bool read_only = read_only_vars_.count(var_id) > 0;
if (auto* tn = storage_type->As<Named>()) {
- read_only = read_only_struct_types_.count(tn->name) > 0;
+ read_only = read_only || read_only_struct_types_.count(tn->name) > 0;
}
// Apply the access(read) or access(read_write) modifier.
@@ -1622,7 +1623,6 @@
const ast::Var* ASTParser::MakeVar(uint32_t id,
core::AddressSpace address_space,
- core::Access access,
const Type* storage_type,
const ast::Expression* initializer,
Attributes attrs) {
@@ -1642,6 +1642,8 @@
return nullptr;
}
+ const auto access = VarAccess(id, storage_type, address_space);
+
// Use type inference if there is an initializer.
auto sym = builder_.Symbols().Register(namer_.Name(id));
return builder_.Var(Source{}, sym, initializer ? ast::Type{} : storage_type->Build(builder_),
@@ -1755,6 +1757,9 @@
}
attrs.Add(builder_.Binding(Source{}, AInt(deco[1])));
}
+ if (deco[0] == uint32_t(spv::Decoration::NonWritable)) {
+ read_only_vars_.insert(id);
+ }
}
if (transfer_pipeline_io) {
diff --git a/src/tint/lang/spirv/reader/ast_parser/ast_parser.h b/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
index 106221a..4174965 100644
--- a/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
+++ b/src/tint/lang/spirv/reader/ast_parser/ast_parser.h
@@ -267,6 +267,7 @@
/// then the `type` parameter is updated. Returns false on failure (with
/// a diagnostic), or when the variable should not be emitted, e.g. for a
/// PointSize builtin.
+ /// This method is idempotent.
/// @param id the ID of the SPIR-V variable
/// @param store_type the WGSL store type for the variable, which should be prepopulated
/// @param attributes the attribute list to populate
@@ -434,17 +435,20 @@
/// @returns a list of SPIR-V decorations.
DecorationList GetMemberPipelineDecorations(const Struct& struct_type, int member_index);
+ /// @param var_id the SPIR-V ID of the OpVariable
/// @param storage_type the 'var' storage type
/// @param address_space the 'var' address space
- /// @returns the access mode for a 'var' declaration with the given storage type and address
- /// space.
- core::Access VarAccess(const Type* storage_type, core::AddressSpace address_space);
+ /// @returns the access mode for a 'var' declaration with the given variable id, storage type
+ /// and address space. Must only be called after decorations for the variable have been
+ /// converted.
+ core::Access VarAccess(uint32_t var_id,
+ const Type* storage_type,
+ core::AddressSpace address_space);
/// Creates an AST 'var' node for a SPIR-V ID, including any attached decorations, unless it's
/// an ignorable builtin variable.
/// @param id the SPIR-V result ID
/// @param address_space the address space, which cannot be core::AddressSpace::kUndefined
- /// @param access the access
/// @param storage_type the storage type of the variable
/// @param initializer the variable initializer
/// @param attributes the variable attributes
@@ -452,7 +456,6 @@
/// in the error case
const ast::Var* MakeVar(uint32_t id,
core::AddressSpace address_space,
- core::Access access,
const Type* storage_type,
const ast::Expression* initializer,
Attributes attributes);
@@ -906,6 +909,9 @@
// The ast::Struct type names with only read-only members.
std::unordered_set<Symbol> read_only_struct_types_;
+ // The IDs of variables marked as NonWritable.
+ std::unordered_set<uint32_t> read_only_vars_;
+
// Maps from OpConstantComposite IDs to identifiers of module-scope const declarations.
std::unordered_map<uint32_t, Symbol> declared_constant_composites_;
diff --git a/src/tint/lang/spirv/reader/ast_parser/function.cc b/src/tint/lang/spirv/reader/ast_parser/function.cc
index b43dc3d..9eaea1e 100644
--- a/src/tint/lang/spirv/reader/ast_parser/function.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/function.cc
@@ -2522,8 +2522,7 @@
}
}
auto* var = parser_impl_.MakeVar(inst.result_id(), core::AddressSpace::kUndefined,
- core::Access::kUndefined, var_store_type, initializer,
- Attributes{});
+ var_store_type, initializer, Attributes{});
auto* var_decl_stmt = create<ast::VariableDeclStatement>(Source{}, var);
AddStatement(var_decl_stmt);
auto* var_type = ty_.Reference(core::AddressSpace::kUndefined, var_store_type);
@@ -3369,9 +3368,8 @@
// no need to remap pointer properties.
auto* store_type = parser_impl_.ConvertType(def_inst->type_id());
AddStatement(create<ast::VariableDeclStatement>(
- Source{},
- parser_impl_.MakeVar(id, core::AddressSpace::kUndefined, core::Access::kUndefined,
- store_type, nullptr, Attributes{})));
+ Source{}, parser_impl_.MakeVar(id, core::AddressSpace::kUndefined, store_type, nullptr,
+ Attributes{})));
auto* type = ty_.Reference(core::AddressSpace::kUndefined, store_type);
identifier_types_.emplace(id, type);
}
diff --git a/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc b/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc
index 3d67797..e3f24e1 100644
--- a/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc
+++ b/src/tint/lang/spirv/reader/ast_parser/module_var_test.cc
@@ -1394,6 +1394,37 @@
<< p->error();
}
+TEST_F(SpvModuleScopeVarParserTest, StorageBuffer_NonWritable_Var) {
+ // Variable should have access(read)
+ auto p = parser(test::Assemble(Preamble() + FragMain() + R"(
+ OpDecorate %s Block
+ OpDecorate %1 DescriptorSet 0
+ OpDecorate %1 Binding 0
+ OpDecorate %1 NonWritable
+ OpMemberDecorate %s 0 Offset 0
+ OpMemberDecorate %s 1 Offset 4
+ %void = OpTypeVoid
+ %voidfn = OpTypeFunction %void
+ %float = OpTypeFloat 32
+
+ %s = OpTypeStruct %float %float
+ %ptr_sb_s = OpTypePointer StorageBuffer %s
+ %1 = OpVariable %ptr_sb_s StorageBuffer
+ )" + MainBody()));
+ ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
+ EXPECT_TRUE(p->error().empty());
+ const auto module_str = test::ToString(p->program());
+ EXPECT_THAT(module_str, HasSubstr(R"(struct S {
+ /* @offset(0) */
+ field0 : f32,
+ /* @offset(4) */
+ field1 : f32,
+}
+
+@group(0) @binding(0) var<storage, read> x_1 : S;
+)")) << module_str;
+}
+
TEST_F(SpvModuleScopeVarParserTest, StorageBuffer_NonWritable_AllMembers) {
// Variable should have access(read)
auto p = parser(test::Assemble(Preamble() + FragMain() + R"(
diff --git a/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl b/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl
index 78faad9..03495e0 100644
--- a/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl
+++ b/test/tint/bug/tint/1776.spvasm.expected.dxc.hlsl
@@ -3,7 +3,7 @@
int b;
};
-RWByteAddressBuffer sb : register(u0);
+ByteAddressBuffer sb : register(t0);
S sb_load(uint offset) {
S tint_symbol = {asfloat(sb.Load4((offset + 0u))), asint(sb.Load((offset + 16u)))};
diff --git a/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl b/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl
index 78faad9..03495e0 100644
--- a/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl
+++ b/test/tint/bug/tint/1776.spvasm.expected.fxc.hlsl
@@ -3,7 +3,7 @@
int b;
};
-RWByteAddressBuffer sb : register(u0);
+ByteAddressBuffer sb : register(t0);
S sb_load(uint offset) {
S tint_symbol = {asfloat(sb.Load4((offset + 0u))), asint(sb.Load((offset + 16u)))};
diff --git a/test/tint/bug/tint/1776.spvasm.expected.msl b/test/tint/bug/tint/1776.spvasm.expected.msl
index 0933ac0..0548772 100644
--- a/test/tint/bug/tint/1776.spvasm.expected.msl
+++ b/test/tint/bug/tint/1776.spvasm.expected.msl
@@ -24,12 +24,12 @@
/* 0x0000 */ tint_array<S, 1> inner;
};
-void main_1(device sb_block* const tint_symbol_1) {
+void main_1(const device sb_block* const tint_symbol_1) {
S const x_18 = (*(tint_symbol_1)).inner[1];
return;
}
-kernel void tint_symbol(device sb_block* tint_symbol_2 [[buffer(0)]]) {
+kernel void tint_symbol(const device sb_block* tint_symbol_2 [[buffer(0)]]) {
main_1(tint_symbol_2);
return;
}
diff --git a/test/tint/bug/tint/1776.spvasm.expected.spvasm b/test/tint/bug/tint/1776.spvasm.expected.spvasm
index 6e076d1..9e220a4 100644
--- a/test/tint/bug/tint/1776.spvasm.expected.spvasm
+++ b/test/tint/bug/tint/1776.spvasm.expected.spvasm
@@ -20,6 +20,7 @@
OpMemberDecorate %S 0 Offset 0
OpMemberDecorate %S 1 Offset 16
OpDecorate %_runtimearr_S ArrayStride 32
+ OpDecorate %sb NonWritable
OpDecorate %sb DescriptorSet 0
OpDecorate %sb Binding 0
%float = OpTypeFloat 32
diff --git a/test/tint/bug/tint/1776.spvasm.expected.wgsl b/test/tint/bug/tint/1776.spvasm.expected.wgsl
index f9988fa..fa35f5e 100644
--- a/test/tint/bug/tint/1776.spvasm.expected.wgsl
+++ b/test/tint/bug/tint/1776.spvasm.expected.wgsl
@@ -12,7 +12,7 @@
inner : RTArr,
}
-@group(0) @binding(0) var<storage, read_write> sb : sb_block;
+@group(0) @binding(0) var<storage, read> sb : sb_block;
fn main_1() {
let x_18 = sb.inner[1i];
