[spirv-reader] Add support for Coherent decoration
Adds support for the decoration in the SPIR-V reader. As Tint's IR has
coherent semantics by default, this is a no-op in the parser.
Change-Id: I416a0c6d623f146cd88c71b15abe5d3bf583ad79
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/248634
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
Auto-Submit: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 5c958a7..d939cf2 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -3434,6 +3434,9 @@
case spv::Decoration::Index:
io_attributes.blend_src = deco->GetSingleWordOperand(2);
break;
+ case spv::Decoration::Coherent:
+ // Tint has coherent memory semantics, so this is a no-op.
+ break;
default:
TINT_UNIMPLEMENTED() << "unhandled decoration " << d;
}
diff --git a/src/tint/lang/spirv/reader/parser/var_test.cc b/src/tint/lang/spirv/reader/parser/var_test.cc
index 7e38025..8618e0c 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -221,6 +221,54 @@
)");
}
+TEST_F(SpirvParserTest, StorageVar_Coherent) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpExtension "SPV_KHR_storage_buffer_storage_class"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %1 "main"
+ OpExecutionMode %1 LocalSize 1 1 1
+ OpDecorate %str Block
+ OpMemberDecorate %str 0 Offset 0
+ OpDecorate %6 DescriptorSet 1
+ OpDecorate %6 Binding 2
+ OpDecorate %6 Coherent
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+ %str = OpTypeStruct %uint
+%_ptr_StorageBuffer_str = OpTypePointer StorageBuffer %str
+ %5 = OpTypeFunction %void
+ %6 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
+ %1 = OpFunction %void None %5
+ %7 = OpLabel
+ %8 = OpAccessChain %_ptr_StorageBuffer_uint %6 %uint_0
+ %9 = OpLoad %uint %8
+ OpStore %8 %9
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+tint_symbol_1 = struct @align(4) {
+ tint_symbol:u32 @offset(0)
+}
+
+$B1: { # root
+ %1:ptr<storage, tint_symbol_1, read_write> = var undef @binding_point(1, 2)
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %3:ptr<storage, u32, read_write> = access %1, 0u
+ %4:u32 = load %3
+ store %3, %4
+ ret
+ }
+}
+)");
+}
+
TEST_F(SpirvParserTest, StorageVar_ReadOnly_And_ReadWrite) {
EXPECT_IR(R"(
OpCapability Shader
