[spirv-reader][ir] Add `SPV_KHR_vulkan_memory_model` as an accepted extension.
Accept the vulkan memory model extension. WGSL is already defined as
conforming to the vulkan memory model, so we can just ignore the
extension.
Bug: 42250952
Change-Id: Ic8ecb836f19acd2e7a5c2811c6da9babea68d360
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/248395
Reviewed-by: James Price <jrprice@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/reader/common/helper_test.h b/src/tint/lang/spirv/reader/common/helper_test.h
index 1d05f18..8c90197 100644
--- a/src/tint/lang/spirv/reader/common/helper_test.h
+++ b/src/tint/lang/spirv/reader/common/helper_test.h
@@ -39,10 +39,12 @@
/// Assemble a textual SPIR-V module into a SPIR-V binary.
/// @param spirv_asm the textual SPIR-V assembly
/// @returns the SPIR-V binary data, or an error string
-inline Result<std::vector<uint32_t>, std::string> Assemble(std::string spirv_asm) {
+inline Result<std::vector<uint32_t>, std::string> Assemble(
+ std::string spirv_asm,
+ spv_target_env spv_version = SPV_ENV_UNIVERSAL_1_0) {
StringStream err;
std::vector<uint32_t> binary;
- spvtools::SpirvTools tools(SPV_ENV_UNIVERSAL_1_0);
+ spvtools::SpirvTools tools(spv_version);
tools.SetMessageConsumer(
[&err](spv_message_level_t, const char*, const spv_position_t& pos, const char* msg) {
err << "SPIR-V assembly failed:" << pos.line << ":" << pos.column << ": " << msg;
diff --git a/src/tint/lang/spirv/reader/parser/helper_test.h b/src/tint/lang/spirv/reader/parser/helper_test.h
index d12b784..6c09bfa 100644
--- a/src/tint/lang/spirv/reader/parser/helper_test.h
+++ b/src/tint/lang/spirv/reader/parser/helper_test.h
@@ -43,14 +43,18 @@
// Helper macro to run the parser and compare the disassembled IR to a string.
// Automatically prefixes the IR disassembly with a newline to improve formatting of tests.
-#define EXPECT_IR(asm, ir) \
+#define EXPECT_IR_SPV(asm, ir, spv_version) \
do { \
- auto result = Run(asm); \
+ auto result = Run(asm, spv_version); \
ASSERT_EQ(result, Success) << result.Failure(); \
auto got = "\n" + result.Get(); \
ASSERT_THAT(got, testing::HasSubstr(ir)) << "GOT:\n" << got << "EXPECTED:\n" << ir; \
} while (false)
+// Helper macro to run the parser and compare the disassembled IR to a string.
+// Automatically prefixes the IR disassembly with a newline to improve formatting of tests.
+#define EXPECT_IR(asm, ir) EXPECT_IR_SPV(asm, ir, SPV_ENV_UNIVERSAL_1_0)
+
/// Base helper class for testing the SPIR-V parser implementation.
template <typename BASE>
class SpirvParserTestHelperBase : public BASE {
@@ -58,9 +62,10 @@
/// Run the parser on a SPIR-V module and return the Tint IR or an error string.
/// @param spirv_asm the SPIR-V assembly to parse
/// @returns the disassembled Tint IR or an error
- Result<std::string> Run(std::string spirv_asm) {
+ Result<std::string> Run(std::string spirv_asm,
+ spv_target_env spv_version = SPV_ENV_UNIVERSAL_1_0) {
// Assemble the SPIR-V input.
- auto binary = Assemble(spirv_asm);
+ auto binary = Assemble(spirv_asm, spv_version);
if (binary != Success) {
return binary.Failure();
}
diff --git a/src/tint/lang/spirv/reader/parser/import_test.cc b/src/tint/lang/spirv/reader/parser/import_test.cc
index 0cfe64a..0f7ac6a 100644
--- a/src/tint/lang/spirv/reader/parser/import_test.cc
+++ b/src/tint/lang/spirv/reader/parser/import_test.cc
@@ -102,6 +102,31 @@
)");
}
+TEST_F(SpirvParserTest, Import_VulkanMemoryModel_IgnoredImport) {
+ EXPECT_IR_SPV(R"(
+ OpCapability Shader
+ OpExtension "SPV_KHR_vulkan_memory_model"
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %void_fn = OpTypeFunction %void
+
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ ret
+ }
+}
+)",
+ SPV_ENV_UNIVERSAL_1_3);
+}
+
// TODO(dsinclair): internal compiler error: TINT_UNIMPLEMENTED unhandled SPIR-V type: [uint32]
TEST_F(SpirvParserTest, DISABLED_Import_NonSemantic_IgnoredExtInsts) {
// This is the clspv-compiled output of this OpenCL C:
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index a0b4f78..b966fa4 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -119,7 +119,11 @@
for (const auto& ext : spirv_context_->extensions()) {
auto name = ext.GetOperand(0).AsString();
if (name != "SPV_KHR_storage_buffer_storage_class" &&
- name != "SPV_KHR_non_semantic_info") {
+ name != "SPV_KHR_non_semantic_info" && //
+ // TODO(423644565): We assume the barriers are correct. We should check for any
+ // operation that makes barrier assumptions that aren't consistent with WGSL and
+ // generate the needed barriers.
+ name != "SPV_KHR_vulkan_memory_model") {
return Failure("SPIR-V extension '" + name + "' is not supported");
}
}
