[spirv-reader] Handle pointer to vector component
Add a new transform that replaces all uses of
pointers-to-vector-components with `load_vector_element` and
`store_vector_element` instructions.
Allow the parser to generate these pointers by enabling the IR
validation capability. Enable two tests that cover this.
Add new tests for the full SPIR-V to core IR flow.
Bug: tint:1907
Change-Id: Ic015cc855b7abb535d23376f88eb4c9f3ba5f086
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/170002
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/api/BUILD.bazel b/src/tint/api/BUILD.bazel
index 32a7773..9f218bd 100644
--- a/src/tint/api/BUILD.bazel
+++ b/src/tint/api/BUILD.bazel
@@ -52,7 +52,6 @@
"//src/tint/lang/core/ir",
"//src/tint/lang/core/type",
"//src/tint/lang/hlsl/writer/common",
- "//src/tint/lang/spirv/reader/common",
"//src/tint/lang/wgsl",
"//src/tint/lang/wgsl/ast",
"//src/tint/lang/wgsl/common",
@@ -93,6 +92,7 @@
}) + select({
":tint_build_spv_reader": [
"//src/tint/lang/spirv/reader",
+ "//src/tint/lang/spirv/reader/common",
],
"//conditions:default": [],
}) + select({
diff --git a/src/tint/api/BUILD.cmake b/src/tint/api/BUILD.cmake
index af664e6..9551324 100644
--- a/src/tint/api/BUILD.cmake
+++ b/src/tint/api/BUILD.cmake
@@ -54,7 +54,6 @@
tint_lang_core_ir
tint_lang_core_type
tint_lang_hlsl_writer_common
- tint_lang_spirv_reader_common
tint_lang_wgsl
tint_lang_wgsl_ast
tint_lang_wgsl_common
@@ -100,6 +99,7 @@
if(TINT_BUILD_SPV_READER)
tint_target_add_dependencies(tint_api lib
tint_lang_spirv_reader
+ tint_lang_spirv_reader_common
)
endif(TINT_BUILD_SPV_READER)
diff --git a/src/tint/api/BUILD.gn b/src/tint/api/BUILD.gn
index ff3f483..80382c0 100644
--- a/src/tint/api/BUILD.gn
+++ b/src/tint/api/BUILD.gn
@@ -51,7 +51,6 @@
"${tint_src_dir}/lang/core/ir",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/lang/hlsl/writer/common",
- "${tint_src_dir}/lang/spirv/reader/common",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl/ast",
"${tint_src_dir}/lang/wgsl/common",
@@ -93,7 +92,10 @@
}
if (tint_build_spv_reader) {
- deps += [ "${tint_src_dir}/lang/spirv/reader" ]
+ deps += [
+ "${tint_src_dir}/lang/spirv/reader",
+ "${tint_src_dir}/lang/spirv/reader/common",
+ ]
}
if (tint_build_spv_writer) {
diff --git a/src/tint/cmd/bench/BUILD.bazel b/src/tint/cmd/bench/BUILD.bazel
index a273bc9..743f284 100644
--- a/src/tint/cmd/bench/BUILD.bazel
+++ b/src/tint/cmd/bench/BUILD.bazel
@@ -49,7 +49,6 @@
"//src/tint/lang/core/constant",
"//src/tint/lang/core/ir",
"//src/tint/lang/core/type",
- "//src/tint/lang/spirv/reader/common",
"//src/tint/lang/wgsl",
"//src/tint/lang/wgsl/ast",
"//src/tint/lang/wgsl/common",
@@ -74,6 +73,7 @@
] + select({
":tint_build_spv_reader": [
"//src/tint/lang/spirv/reader",
+ "//src/tint/lang/spirv/reader/common",
],
"//conditions:default": [],
}) + select({
diff --git a/src/tint/cmd/bench/BUILD.cmake b/src/tint/cmd/bench/BUILD.cmake
index ee4b122..cdc5971 100644
--- a/src/tint/cmd/bench/BUILD.cmake
+++ b/src/tint/cmd/bench/BUILD.cmake
@@ -126,7 +126,6 @@
tint_lang_core_constant
tint_lang_core_ir
tint_lang_core_type
- tint_lang_spirv_reader_common
tint_lang_wgsl
tint_lang_wgsl_ast
tint_lang_wgsl_common
@@ -156,6 +155,7 @@
if(TINT_BUILD_SPV_READER)
tint_target_add_dependencies(tint_cmd_bench_bench bench
tint_lang_spirv_reader
+ tint_lang_spirv_reader_common
)
endif(TINT_BUILD_SPV_READER)
diff --git a/src/tint/cmd/bench/BUILD.gn b/src/tint/cmd/bench/BUILD.gn
index 7eb287d..38bdc6a 100644
--- a/src/tint/cmd/bench/BUILD.gn
+++ b/src/tint/cmd/bench/BUILD.gn
@@ -54,7 +54,6 @@
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/ir",
"${tint_src_dir}/lang/core/type",
- "${tint_src_dir}/lang/spirv/reader/common",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl/ast",
"${tint_src_dir}/lang/wgsl/common",
@@ -78,7 +77,10 @@
]
if (tint_build_spv_reader) {
- deps += [ "${tint_src_dir}/lang/spirv/reader" ]
+ deps += [
+ "${tint_src_dir}/lang/spirv/reader",
+ "${tint_src_dir}/lang/spirv/reader/common",
+ ]
}
if (tint_build_wgsl_reader) {
diff --git a/src/tint/cmd/common/BUILD.bazel b/src/tint/cmd/common/BUILD.bazel
index c768f73..affb277 100644
--- a/src/tint/cmd/common/BUILD.bazel
+++ b/src/tint/cmd/common/BUILD.bazel
@@ -53,7 +53,6 @@
"//src/tint/lang/core/constant",
"//src/tint/lang/core/ir",
"//src/tint/lang/core/type",
- "//src/tint/lang/spirv/reader/common",
"//src/tint/lang/wgsl",
"//src/tint/lang/wgsl/ast",
"//src/tint/lang/wgsl/common",
@@ -78,6 +77,7 @@
] + select({
":tint_build_spv_reader": [
"//src/tint/lang/spirv/reader",
+ "//src/tint/lang/spirv/reader/common",
],
"//conditions:default": [],
}) + select({
diff --git a/src/tint/cmd/common/BUILD.cmake b/src/tint/cmd/common/BUILD.cmake
index 43d472f..2c0d823 100644
--- a/src/tint/cmd/common/BUILD.cmake
+++ b/src/tint/cmd/common/BUILD.cmake
@@ -52,7 +52,6 @@
tint_lang_core_constant
tint_lang_core_ir
tint_lang_core_type
- tint_lang_spirv_reader_common
tint_lang_wgsl
tint_lang_wgsl_ast
tint_lang_wgsl_common
@@ -79,6 +78,7 @@
if(TINT_BUILD_SPV_READER)
tint_target_add_dependencies(tint_cmd_common lib
tint_lang_spirv_reader
+ tint_lang_spirv_reader_common
)
endif(TINT_BUILD_SPV_READER)
diff --git a/src/tint/cmd/common/BUILD.gn b/src/tint/cmd/common/BUILD.gn
index 4b1f5e3..6026bf5 100644
--- a/src/tint/cmd/common/BUILD.gn
+++ b/src/tint/cmd/common/BUILD.gn
@@ -56,7 +56,6 @@
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/ir",
"${tint_src_dir}/lang/core/type",
- "${tint_src_dir}/lang/spirv/reader/common",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl/ast",
"${tint_src_dir}/lang/wgsl/common",
@@ -81,7 +80,10 @@
]
if (tint_build_spv_reader) {
- deps += [ "${tint_src_dir}/lang/spirv/reader" ]
+ deps += [
+ "${tint_src_dir}/lang/spirv/reader",
+ "${tint_src_dir}/lang/spirv/reader/common",
+ ]
}
if (tint_build_spv_reader || tint_build_spv_writer) {
diff --git a/src/tint/cmd/info/BUILD.bazel b/src/tint/cmd/info/BUILD.bazel
index c19a5da..6d0b8f9 100644
--- a/src/tint/cmd/info/BUILD.bazel
+++ b/src/tint/cmd/info/BUILD.bazel
@@ -47,7 +47,6 @@
"//src/tint/lang/core",
"//src/tint/lang/core/constant",
"//src/tint/lang/core/type",
- "//src/tint/lang/spirv/reader/common",
"//src/tint/lang/wgsl",
"//src/tint/lang/wgsl/ast",
"//src/tint/lang/wgsl/common",
@@ -70,6 +69,11 @@
"//src/tint/utils/text",
"//src/tint/utils/traits",
] + select({
+ ":tint_build_spv_reader": [
+ "//src/tint/lang/spirv/reader/common",
+ ],
+ "//conditions:default": [],
+ }) + select({
":tint_build_spv_reader_or_tint_build_spv_writer": [
"@spirv_tools",
],
diff --git a/src/tint/cmd/info/BUILD.cmake b/src/tint/cmd/info/BUILD.cmake
index 74d28e0..dda56b9 100644
--- a/src/tint/cmd/info/BUILD.cmake
+++ b/src/tint/cmd/info/BUILD.cmake
@@ -48,7 +48,6 @@
tint_lang_core
tint_lang_core_constant
tint_lang_core_type
- tint_lang_spirv_reader_common
tint_lang_wgsl
tint_lang_wgsl_ast
tint_lang_wgsl_common
@@ -72,6 +71,12 @@
tint_utils_traits
)
+if(TINT_BUILD_SPV_READER)
+ tint_target_add_dependencies(tint_cmd_info_cmd cmd
+ tint_lang_spirv_reader_common
+ )
+endif(TINT_BUILD_SPV_READER)
+
if(TINT_BUILD_SPV_READER OR TINT_BUILD_SPV_WRITER)
tint_target_add_external_dependencies(tint_cmd_info_cmd cmd
"spirv-tools"
diff --git a/src/tint/cmd/info/BUILD.gn b/src/tint/cmd/info/BUILD.gn
index 6f3a945..5f06fcd 100644
--- a/src/tint/cmd/info/BUILD.gn
+++ b/src/tint/cmd/info/BUILD.gn
@@ -47,7 +47,6 @@
"${tint_src_dir}/lang/core",
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/type",
- "${tint_src_dir}/lang/spirv/reader/common",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl/ast",
"${tint_src_dir}/lang/wgsl/common",
@@ -71,6 +70,10 @@
"${tint_src_dir}/utils/traits",
]
+ if (tint_build_spv_reader) {
+ deps += [ "${tint_src_dir}/lang/spirv/reader/common" ]
+ }
+
if (tint_build_spv_reader || tint_build_spv_writer) {
deps += [
"${tint_spirv_tools_dir}:spvtools_headers",
diff --git a/src/tint/cmd/loopy/BUILD.bazel b/src/tint/cmd/loopy/BUILD.bazel
index e0b4ebe..1b0859f 100644
--- a/src/tint/cmd/loopy/BUILD.bazel
+++ b/src/tint/cmd/loopy/BUILD.bazel
@@ -51,7 +51,6 @@
"//src/tint/lang/core/ir",
"//src/tint/lang/core/type",
"//src/tint/lang/hlsl/writer/common",
- "//src/tint/lang/spirv/reader/common",
"//src/tint/lang/wgsl",
"//src/tint/lang/wgsl/ast",
"//src/tint/lang/wgsl/common",
@@ -95,6 +94,7 @@
}) + select({
":tint_build_spv_reader": [
"//src/tint/lang/spirv/reader",
+ "//src/tint/lang/spirv/reader/common",
],
"//conditions:default": [],
}) + select({
diff --git a/src/tint/cmd/loopy/BUILD.cmake b/src/tint/cmd/loopy/BUILD.cmake
index abf2e07..2597a47 100644
--- a/src/tint/cmd/loopy/BUILD.cmake
+++ b/src/tint/cmd/loopy/BUILD.cmake
@@ -52,7 +52,6 @@
tint_lang_core_ir
tint_lang_core_type
tint_lang_hlsl_writer_common
- tint_lang_spirv_reader_common
tint_lang_wgsl
tint_lang_wgsl_ast
tint_lang_wgsl_common
@@ -101,6 +100,7 @@
if(TINT_BUILD_SPV_READER)
tint_target_add_dependencies(tint_cmd_loopy_cmd cmd
tint_lang_spirv_reader
+ tint_lang_spirv_reader_common
)
endif(TINT_BUILD_SPV_READER)
diff --git a/src/tint/cmd/loopy/BUILD.gn b/src/tint/cmd/loopy/BUILD.gn
index 3241429..5176168 100644
--- a/src/tint/cmd/loopy/BUILD.gn
+++ b/src/tint/cmd/loopy/BUILD.gn
@@ -51,7 +51,6 @@
"${tint_src_dir}/lang/core/ir",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/lang/hlsl/writer/common",
- "${tint_src_dir}/lang/spirv/reader/common",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl/ast",
"${tint_src_dir}/lang/wgsl/common",
@@ -96,7 +95,10 @@
}
if (tint_build_spv_reader) {
- deps += [ "${tint_src_dir}/lang/spirv/reader" ]
+ deps += [
+ "${tint_src_dir}/lang/spirv/reader",
+ "${tint_src_dir}/lang/spirv/reader/common",
+ ]
}
if (tint_build_spv_writer) {
diff --git a/src/tint/cmd/test/BUILD.bazel b/src/tint/cmd/test/BUILD.bazel
index a44e992..01d6cf6 100644
--- a/src/tint/cmd/test/BUILD.bazel
+++ b/src/tint/cmd/test/BUILD.bazel
@@ -52,6 +52,7 @@
"//src/tint/lang/core:test",
"//src/tint/lang/msl/ir:test",
"//src/tint/lang/spirv/ir:test",
+ "//src/tint/lang/spirv/reader/lower:test",
"//src/tint/lang/wgsl/ast:test",
"//src/tint/lang/wgsl/helpers:test",
"//src/tint/lang/wgsl/program:test",
@@ -122,7 +123,9 @@
"//conditions:default": [],
}) + select({
":tint_build_spv_reader": [
+ "//src/tint/lang/spirv/reader/common:test",
"//src/tint/lang/spirv/reader/parser:test",
+ "//src/tint/lang/spirv/reader:test",
],
"//conditions:default": [],
}) + select({
diff --git a/src/tint/cmd/test/BUILD.cmake b/src/tint/cmd/test/BUILD.cmake
index e79c221..d81fdcb 100644
--- a/src/tint/cmd/test/BUILD.cmake
+++ b/src/tint/cmd/test/BUILD.cmake
@@ -53,6 +53,7 @@
tint_lang_core_test
tint_lang_msl_ir_test
tint_lang_spirv_ir_test
+ tint_lang_spirv_reader_lower_test
tint_lang_wgsl_ast_test
tint_lang_wgsl_helpers_test
tint_lang_wgsl_program_test
@@ -135,7 +136,9 @@
if(TINT_BUILD_SPV_READER)
tint_target_add_dependencies(tint_cmd_test_test_cmd test_cmd
+ tint_lang_spirv_reader_common_test
tint_lang_spirv_reader_parser_test
+ tint_lang_spirv_reader_test
)
endif(TINT_BUILD_SPV_READER)
diff --git a/src/tint/cmd/test/BUILD.gn b/src/tint/cmd/test/BUILD.gn
index 55c146b..abb1fd2 100644
--- a/src/tint/cmd/test/BUILD.gn
+++ b/src/tint/cmd/test/BUILD.gn
@@ -58,6 +58,7 @@
"${tint_src_dir}/lang/core/type:unittests",
"${tint_src_dir}/lang/msl/ir:unittests",
"${tint_src_dir}/lang/spirv/ir:unittests",
+ "${tint_src_dir}/lang/spirv/reader/lower:unittests",
"${tint_src_dir}/lang/wgsl:unittests",
"${tint_src_dir}/lang/wgsl/ast:unittests",
"${tint_src_dir}/lang/wgsl/helpers:unittests",
@@ -128,7 +129,11 @@
}
if (tint_build_spv_reader) {
- deps += [ "${tint_src_dir}/lang/spirv/reader/parser:unittests" ]
+ deps += [
+ "${tint_src_dir}/lang/spirv/reader:unittests",
+ "${tint_src_dir}/lang/spirv/reader/common:unittests",
+ "${tint_src_dir}/lang/spirv/reader/parser:unittests",
+ ]
}
if (tint_build_spv_reader && tint_build_wgsl_reader &&
diff --git a/src/tint/cmd/tint/BUILD.bazel b/src/tint/cmd/tint/BUILD.bazel
index 934ff47..fbc1286 100644
--- a/src/tint/cmd/tint/BUILD.bazel
+++ b/src/tint/cmd/tint/BUILD.bazel
@@ -51,7 +51,6 @@
"//src/tint/lang/core/ir",
"//src/tint/lang/core/type",
"//src/tint/lang/hlsl/writer/common",
- "//src/tint/lang/spirv/reader/common",
"//src/tint/lang/wgsl",
"//src/tint/lang/wgsl/ast",
"//src/tint/lang/wgsl/ast/transform",
@@ -104,6 +103,11 @@
],
"//conditions:default": [],
}) + select({
+ ":tint_build_spv_reader": [
+ "//src/tint/lang/spirv/reader/common",
+ ],
+ "//conditions:default": [],
+ }) + select({
":tint_build_spv_reader_or_tint_build_spv_writer": [
"@spirv_tools",
],
diff --git a/src/tint/cmd/tint/BUILD.cmake b/src/tint/cmd/tint/BUILD.cmake
index 3924e7d..9a02dfe 100644
--- a/src/tint/cmd/tint/BUILD.cmake
+++ b/src/tint/cmd/tint/BUILD.cmake
@@ -52,7 +52,6 @@
tint_lang_core_ir
tint_lang_core_type
tint_lang_hlsl_writer_common
- tint_lang_spirv_reader_common
tint_lang_wgsl
tint_lang_wgsl_ast
tint_lang_wgsl_ast_transform
@@ -110,6 +109,12 @@
)
endif(TINT_BUILD_MSL_WRITER)
+if(TINT_BUILD_SPV_READER)
+ tint_target_add_dependencies(tint_cmd_tint_cmd cmd
+ tint_lang_spirv_reader_common
+ )
+endif(TINT_BUILD_SPV_READER)
+
if(TINT_BUILD_SPV_READER OR TINT_BUILD_SPV_WRITER)
tint_target_add_external_dependencies(tint_cmd_tint_cmd cmd
"spirv-tools"
diff --git a/src/tint/cmd/tint/BUILD.gn b/src/tint/cmd/tint/BUILD.gn
index b9fca30..5d0447f 100644
--- a/src/tint/cmd/tint/BUILD.gn
+++ b/src/tint/cmd/tint/BUILD.gn
@@ -51,7 +51,6 @@
"${tint_src_dir}/lang/core/ir",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/lang/hlsl/writer/common",
- "${tint_src_dir}/lang/spirv/reader/common",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl/ast",
"${tint_src_dir}/lang/wgsl/ast/transform",
@@ -107,6 +106,10 @@
]
}
+ if (tint_build_spv_reader) {
+ deps += [ "${tint_src_dir}/lang/spirv/reader/common" ]
+ }
+
if (tint_build_spv_reader || tint_build_spv_writer) {
deps += [
"${tint_spirv_tools_dir}:spvtools_headers",
diff --git a/src/tint/lang/spirv/reader/BUILD.bazel b/src/tint/lang/spirv/reader/BUILD.bazel
index f7662cf..3f50893 100644
--- a/src/tint/lang/spirv/reader/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/BUILD.bazel
@@ -50,7 +50,7 @@
"//src/tint/lang/core/constant",
"//src/tint/lang/core/ir",
"//src/tint/lang/core/type",
- "//src/tint/lang/spirv/reader/common",
+ "//src/tint/lang/spirv/reader/lower",
"//src/tint/lang/wgsl",
"//src/tint/lang/wgsl/ast",
"//src/tint/lang/wgsl/common",
@@ -73,6 +73,7 @@
] + select({
":tint_build_spv_reader": [
"//src/tint/lang/spirv/reader/ast_parser",
+ "//src/tint/lang/spirv/reader/common",
"//src/tint/lang/spirv/reader/parser",
],
"//conditions:default": [],
@@ -80,9 +81,70 @@
copts = COPTS,
visibility = ["//visibility:public"],
)
+cc_library(
+ name = "test",
+ alwayslink = True,
+ srcs = [
+ "reader_test.cc",
+ ],
+ deps = [
+ "//src/tint/api/common",
+ "//src/tint/lang/core",
+ "//src/tint/lang/core/constant",
+ "//src/tint/lang/core/ir",
+ "//src/tint/lang/core/type",
+ "//src/tint/lang/wgsl",
+ "//src/tint/lang/wgsl/ast",
+ "//src/tint/lang/wgsl/common",
+ "//src/tint/lang/wgsl/features",
+ "//src/tint/lang/wgsl/program",
+ "//src/tint/lang/wgsl/sem",
+ "//src/tint/utils/containers",
+ "//src/tint/utils/diagnostic",
+ "//src/tint/utils/ice",
+ "//src/tint/utils/id",
+ "//src/tint/utils/macros",
+ "//src/tint/utils/math",
+ "//src/tint/utils/memory",
+ "//src/tint/utils/reflection",
+ "//src/tint/utils/result",
+ "//src/tint/utils/rtti",
+ "//src/tint/utils/symbol",
+ "//src/tint/utils/text",
+ "//src/tint/utils/traits",
+ "@gtest",
+ ] + select({
+ ":tint_build_spv_reader": [
+ "//src/tint/lang/spirv/reader",
+ "//src/tint/lang/spirv/reader/common",
+ "//src/tint/lang/spirv/reader/common:test",
+ ],
+ "//conditions:default": [],
+ }) + select({
+ ":tint_build_spv_reader_or_tint_build_spv_writer": [
+ "@spirv_tools",
+ ],
+ "//conditions:default": [],
+ }),
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
alias(
name = "tint_build_spv_reader",
actual = "//src/tint:tint_build_spv_reader_true",
)
+alias(
+ name = "tint_build_spv_writer",
+ actual = "//src/tint:tint_build_spv_writer_true",
+)
+
+selects.config_setting_group(
+ name = "tint_build_spv_reader_or_tint_build_spv_writer",
+ match_any = [
+ "tint_build_spv_reader",
+ "tint_build_spv_writer",
+ ],
+)
+
diff --git a/src/tint/lang/spirv/reader/BUILD.cmake b/src/tint/lang/spirv/reader/BUILD.cmake
index 6e0d5ac..6b1a73c 100644
--- a/src/tint/lang/spirv/reader/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/BUILD.cmake
@@ -37,6 +37,7 @@
include(lang/spirv/reader/ast_lower/BUILD.cmake)
include(lang/spirv/reader/ast_parser/BUILD.cmake)
include(lang/spirv/reader/common/BUILD.cmake)
+include(lang/spirv/reader/lower/BUILD.cmake)
include(lang/spirv/reader/parser/BUILD.cmake)
if(TINT_BUILD_SPV_READER)
@@ -56,7 +57,7 @@
tint_lang_core_constant
tint_lang_core_ir
tint_lang_core_type
- tint_lang_spirv_reader_common
+ tint_lang_spirv_reader_lower
tint_lang_wgsl
tint_lang_wgsl_ast
tint_lang_wgsl_common
@@ -81,8 +82,65 @@
if(TINT_BUILD_SPV_READER)
tint_target_add_dependencies(tint_lang_spirv_reader lib
tint_lang_spirv_reader_ast_parser
+ tint_lang_spirv_reader_common
tint_lang_spirv_reader_parser
)
endif(TINT_BUILD_SPV_READER)
+endif(TINT_BUILD_SPV_READER)
+if(TINT_BUILD_SPV_READER)
+################################################################################
+# Target: tint_lang_spirv_reader_test
+# Kind: test
+# Condition: TINT_BUILD_SPV_READER
+################################################################################
+tint_add_target(tint_lang_spirv_reader_test test
+ lang/spirv/reader/reader_test.cc
+)
+
+tint_target_add_dependencies(tint_lang_spirv_reader_test test
+ tint_api_common
+ tint_lang_core
+ tint_lang_core_constant
+ tint_lang_core_ir
+ tint_lang_core_type
+ tint_lang_wgsl
+ tint_lang_wgsl_ast
+ tint_lang_wgsl_common
+ tint_lang_wgsl_features
+ tint_lang_wgsl_program
+ tint_lang_wgsl_sem
+ tint_utils_containers
+ tint_utils_diagnostic
+ tint_utils_ice
+ tint_utils_id
+ tint_utils_macros
+ tint_utils_math
+ tint_utils_memory
+ tint_utils_reflection
+ tint_utils_result
+ tint_utils_rtti
+ tint_utils_symbol
+ tint_utils_text
+ tint_utils_traits
+)
+
+tint_target_add_external_dependencies(tint_lang_spirv_reader_test test
+ "gtest"
+)
+
+if(TINT_BUILD_SPV_READER)
+ tint_target_add_dependencies(tint_lang_spirv_reader_test test
+ tint_lang_spirv_reader
+ tint_lang_spirv_reader_common
+ tint_lang_spirv_reader_common_test
+ )
+endif(TINT_BUILD_SPV_READER)
+
+if(TINT_BUILD_SPV_READER OR TINT_BUILD_SPV_WRITER)
+ tint_target_add_external_dependencies(tint_lang_spirv_reader_test test
+ "spirv-tools"
+ )
+endif(TINT_BUILD_SPV_READER OR TINT_BUILD_SPV_WRITER)
+
endif(TINT_BUILD_SPV_READER)
\ No newline at end of file
diff --git a/src/tint/lang/spirv/reader/BUILD.gn b/src/tint/lang/spirv/reader/BUILD.gn
index 82b4314..09788fc 100644
--- a/src/tint/lang/spirv/reader/BUILD.gn
+++ b/src/tint/lang/spirv/reader/BUILD.gn
@@ -37,6 +37,10 @@
import("../../../../../scripts/tint_overrides_with_defaults.gni")
import("${tint_src_dir}/tint.gni")
+
+if (tint_build_unittests || tint_build_benchmarks) {
+ import("//testing/test.gni")
+}
if (tint_build_spv_reader) {
libtint_source_set("reader") {
sources = [
@@ -49,7 +53,7 @@
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/ir",
"${tint_src_dir}/lang/core/type",
- "${tint_src_dir}/lang/spirv/reader/common",
+ "${tint_src_dir}/lang/spirv/reader/lower",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl/ast",
"${tint_src_dir}/lang/wgsl/common",
@@ -74,8 +78,58 @@
if (tint_build_spv_reader) {
deps += [
"${tint_src_dir}/lang/spirv/reader/ast_parser",
+ "${tint_src_dir}/lang/spirv/reader/common",
"${tint_src_dir}/lang/spirv/reader/parser",
]
}
}
}
+if (tint_build_unittests) {
+ if (tint_build_spv_reader) {
+ tint_unittests_source_set("unittests") {
+ sources = [ "reader_test.cc" ]
+ deps = [
+ "${tint_src_dir}:gmock_and_gtest",
+ "${tint_src_dir}/api/common",
+ "${tint_src_dir}/lang/core",
+ "${tint_src_dir}/lang/core/constant",
+ "${tint_src_dir}/lang/core/ir",
+ "${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/lang/wgsl",
+ "${tint_src_dir}/lang/wgsl/ast",
+ "${tint_src_dir}/lang/wgsl/common",
+ "${tint_src_dir}/lang/wgsl/features",
+ "${tint_src_dir}/lang/wgsl/program",
+ "${tint_src_dir}/lang/wgsl/sem",
+ "${tint_src_dir}/utils/containers",
+ "${tint_src_dir}/utils/diagnostic",
+ "${tint_src_dir}/utils/ice",
+ "${tint_src_dir}/utils/id",
+ "${tint_src_dir}/utils/macros",
+ "${tint_src_dir}/utils/math",
+ "${tint_src_dir}/utils/memory",
+ "${tint_src_dir}/utils/reflection",
+ "${tint_src_dir}/utils/result",
+ "${tint_src_dir}/utils/rtti",
+ "${tint_src_dir}/utils/symbol",
+ "${tint_src_dir}/utils/text",
+ "${tint_src_dir}/utils/traits",
+ ]
+
+ if (tint_build_spv_reader) {
+ deps += [
+ "${tint_src_dir}/lang/spirv/reader",
+ "${tint_src_dir}/lang/spirv/reader/common",
+ "${tint_src_dir}/lang/spirv/reader/common:unittests",
+ ]
+ }
+
+ if (tint_build_spv_reader || tint_build_spv_writer) {
+ deps += [
+ "${tint_spirv_tools_dir}:spvtools_headers",
+ "${tint_spirv_tools_dir}:spvtools_val",
+ ]
+ }
+ }
+ }
+}
diff --git a/src/tint/lang/spirv/reader/ast_parser/BUILD.bazel b/src/tint/lang/spirv/reader/ast_parser/BUILD.bazel
index 421ec97..c656285 100644
--- a/src/tint/lang/spirv/reader/ast_parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/ast_parser/BUILD.bazel
@@ -67,7 +67,6 @@
"//src/tint/lang/core",
"//src/tint/lang/core/constant",
"//src/tint/lang/core/type",
- "//src/tint/lang/spirv/reader/common",
"//src/tint/lang/wgsl",
"//src/tint/lang/wgsl/ast",
"//src/tint/lang/wgsl/ast/transform",
@@ -92,6 +91,7 @@
] + select({
":tint_build_spv_reader": [
"//src/tint/lang/spirv/reader/ast_lower",
+ "//src/tint/lang/spirv/reader/common",
],
"//conditions:default": [],
}) + select({
@@ -148,7 +148,6 @@
"//src/tint/lang/core",
"//src/tint/lang/core/constant",
"//src/tint/lang/core/type",
- "//src/tint/lang/spirv/reader/common",
"//src/tint/lang/wgsl",
"//src/tint/lang/wgsl/ast",
"//src/tint/lang/wgsl/common",
@@ -173,6 +172,7 @@
] + select({
":tint_build_spv_reader": [
"//src/tint/lang/spirv/reader/ast_parser",
+ "//src/tint/lang/spirv/reader/common",
],
"//conditions:default": [],
}) + select({
diff --git a/src/tint/lang/spirv/reader/ast_parser/BUILD.cmake b/src/tint/lang/spirv/reader/ast_parser/BUILD.cmake
index cf9191f..e4ab32d 100644
--- a/src/tint/lang/spirv/reader/ast_parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/ast_parser/BUILD.cmake
@@ -68,7 +68,6 @@
tint_lang_core
tint_lang_core_constant
tint_lang_core_type
- tint_lang_spirv_reader_common
tint_lang_wgsl
tint_lang_wgsl_ast
tint_lang_wgsl_ast_transform
@@ -95,6 +94,7 @@
if(TINT_BUILD_SPV_READER)
tint_target_add_dependencies(tint_lang_spirv_reader_ast_parser lib
tint_lang_spirv_reader_ast_lower
+ tint_lang_spirv_reader_common
)
endif(TINT_BUILD_SPV_READER)
@@ -154,7 +154,6 @@
tint_lang_core
tint_lang_core_constant
tint_lang_core_type
- tint_lang_spirv_reader_common
tint_lang_wgsl
tint_lang_wgsl_ast
tint_lang_wgsl_common
@@ -184,6 +183,7 @@
if(TINT_BUILD_SPV_READER)
tint_target_add_dependencies(tint_lang_spirv_reader_ast_parser_test test
tint_lang_spirv_reader_ast_parser
+ tint_lang_spirv_reader_common
)
endif(TINT_BUILD_SPV_READER)
diff --git a/src/tint/lang/spirv/reader/ast_parser/BUILD.gn b/src/tint/lang/spirv/reader/ast_parser/BUILD.gn
index f82d2fb..5c8e0c1 100644
--- a/src/tint/lang/spirv/reader/ast_parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/ast_parser/BUILD.gn
@@ -70,7 +70,6 @@
"${tint_src_dir}/lang/core",
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/type",
- "${tint_src_dir}/lang/spirv/reader/common",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl/ast",
"${tint_src_dir}/lang/wgsl/ast/transform",
@@ -95,7 +94,10 @@
]
if (tint_build_spv_reader) {
- deps += [ "${tint_src_dir}/lang/spirv/reader/ast_lower" ]
+ deps += [
+ "${tint_src_dir}/lang/spirv/reader/ast_lower",
+ "${tint_src_dir}/lang/spirv/reader/common",
+ ]
}
if (tint_build_spv_reader || tint_build_spv_writer) {
@@ -154,7 +156,6 @@
"${tint_src_dir}/lang/core",
"${tint_src_dir}/lang/core/constant",
"${tint_src_dir}/lang/core/type",
- "${tint_src_dir}/lang/spirv/reader/common",
"${tint_src_dir}/lang/wgsl",
"${tint_src_dir}/lang/wgsl/ast",
"${tint_src_dir}/lang/wgsl/common",
@@ -178,7 +179,10 @@
]
if (tint_build_spv_reader) {
- deps += [ "${tint_src_dir}/lang/spirv/reader/ast_parser" ]
+ deps += [
+ "${tint_src_dir}/lang/spirv/reader/ast_parser",
+ "${tint_src_dir}/lang/spirv/reader/common",
+ ]
}
if (tint_build_spv_reader || tint_build_spv_writer) {
diff --git a/src/tint/lang/spirv/reader/common/BUILD.bazel b/src/tint/lang/spirv/reader/common/BUILD.bazel
index b74c742..438bddd 100644
--- a/src/tint/lang/spirv/reader/common/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/common/BUILD.bazel
@@ -60,4 +60,48 @@
copts = COPTS,
visibility = ["//visibility:public"],
)
+cc_library(
+ name = "test",
+ alwayslink = True,
+ srcs = [
+ "helper_test.h",
+ ],
+ deps = [
+ "//src/tint/utils/containers",
+ "//src/tint/utils/diagnostic",
+ "//src/tint/utils/ice",
+ "//src/tint/utils/macros",
+ "//src/tint/utils/math",
+ "//src/tint/utils/memory",
+ "//src/tint/utils/result",
+ "//src/tint/utils/rtti",
+ "//src/tint/utils/text",
+ "//src/tint/utils/traits",
+ ] + select({
+ ":tint_build_spv_reader_or_tint_build_spv_writer": [
+ "@spirv_tools",
+ ],
+ "//conditions:default": [],
+ }),
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
+
+alias(
+ name = "tint_build_spv_reader",
+ actual = "//src/tint:tint_build_spv_reader_true",
+)
+
+alias(
+ name = "tint_build_spv_writer",
+ actual = "//src/tint:tint_build_spv_writer_true",
+)
+
+selects.config_setting_group(
+ name = "tint_build_spv_reader_or_tint_build_spv_writer",
+ match_any = [
+ "tint_build_spv_reader",
+ "tint_build_spv_writer",
+ ],
+)
diff --git a/src/tint/lang/spirv/reader/common/BUILD.cfg b/src/tint/lang/spirv/reader/common/BUILD.cfg
new file mode 100644
index 0000000..a460fd5
--- /dev/null
+++ b/src/tint/lang/spirv/reader/common/BUILD.cfg
@@ -0,0 +1,3 @@
+{
+ "condition": "tint_build_spv_reader"
+}
diff --git a/src/tint/lang/spirv/reader/common/BUILD.cmake b/src/tint/lang/spirv/reader/common/BUILD.cmake
index 23c0ff1..f0206ba 100644
--- a/src/tint/lang/spirv/reader/common/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/common/BUILD.cmake
@@ -34,9 +34,11 @@
# Do not modify this file directly
################################################################################
+if(TINT_BUILD_SPV_READER)
################################################################################
# Target: tint_lang_spirv_reader_common
# Kind: lib
+# Condition: TINT_BUILD_SPV_READER
################################################################################
tint_add_target(tint_lang_spirv_reader_common lib
lang/spirv/reader/common/common.cc
@@ -56,3 +58,35 @@
tint_utils_rtti
tint_utils_traits
)
+
+endif(TINT_BUILD_SPV_READER)
+if(TINT_BUILD_SPV_READER)
+################################################################################
+# Target: tint_lang_spirv_reader_common_test
+# Kind: test
+# Condition: TINT_BUILD_SPV_READER
+################################################################################
+tint_add_target(tint_lang_spirv_reader_common_test test
+ lang/spirv/reader/common/helper_test.h
+)
+
+tint_target_add_dependencies(tint_lang_spirv_reader_common_test test
+ tint_utils_containers
+ tint_utils_diagnostic
+ tint_utils_ice
+ tint_utils_macros
+ tint_utils_math
+ tint_utils_memory
+ tint_utils_result
+ tint_utils_rtti
+ tint_utils_text
+ tint_utils_traits
+)
+
+if(TINT_BUILD_SPV_READER OR TINT_BUILD_SPV_WRITER)
+ tint_target_add_external_dependencies(tint_lang_spirv_reader_common_test test
+ "spirv-tools"
+ )
+endif(TINT_BUILD_SPV_READER OR TINT_BUILD_SPV_WRITER)
+
+endif(TINT_BUILD_SPV_READER)
\ No newline at end of file
diff --git a/src/tint/lang/spirv/reader/common/BUILD.gn b/src/tint/lang/spirv/reader/common/BUILD.gn
index 9e75c83..8383ffe 100644
--- a/src/tint/lang/spirv/reader/common/BUILD.gn
+++ b/src/tint/lang/spirv/reader/common/BUILD.gn
@@ -38,22 +38,53 @@
import("${tint_src_dir}/tint.gni")
-libtint_source_set("common") {
- sources = [
- "common.cc",
- "options.h",
- ]
- deps = [
- "${tint_src_dir}/lang/wgsl",
- "${tint_src_dir}/lang/wgsl/common",
- "${tint_src_dir}/lang/wgsl/features",
- "${tint_src_dir}/utils/containers",
- "${tint_src_dir}/utils/ice",
- "${tint_src_dir}/utils/macros",
- "${tint_src_dir}/utils/math",
- "${tint_src_dir}/utils/memory",
- "${tint_src_dir}/utils/reflection",
- "${tint_src_dir}/utils/rtti",
- "${tint_src_dir}/utils/traits",
- ]
+if (tint_build_unittests || tint_build_benchmarks) {
+ import("//testing/test.gni")
+}
+if (tint_build_spv_reader) {
+ libtint_source_set("common") {
+ sources = [
+ "common.cc",
+ "options.h",
+ ]
+ deps = [
+ "${tint_src_dir}/lang/wgsl",
+ "${tint_src_dir}/lang/wgsl/common",
+ "${tint_src_dir}/lang/wgsl/features",
+ "${tint_src_dir}/utils/containers",
+ "${tint_src_dir}/utils/ice",
+ "${tint_src_dir}/utils/macros",
+ "${tint_src_dir}/utils/math",
+ "${tint_src_dir}/utils/memory",
+ "${tint_src_dir}/utils/reflection",
+ "${tint_src_dir}/utils/rtti",
+ "${tint_src_dir}/utils/traits",
+ ]
+ }
+}
+if (tint_build_unittests) {
+ if (tint_build_spv_reader) {
+ tint_unittests_source_set("unittests") {
+ sources = [ "helper_test.h" ]
+ deps = [
+ "${tint_src_dir}/utils/containers",
+ "${tint_src_dir}/utils/diagnostic",
+ "${tint_src_dir}/utils/ice",
+ "${tint_src_dir}/utils/macros",
+ "${tint_src_dir}/utils/math",
+ "${tint_src_dir}/utils/memory",
+ "${tint_src_dir}/utils/result",
+ "${tint_src_dir}/utils/rtti",
+ "${tint_src_dir}/utils/text",
+ "${tint_src_dir}/utils/traits",
+ ]
+
+ if (tint_build_spv_reader || tint_build_spv_writer) {
+ deps += [
+ "${tint_spirv_tools_dir}:spvtools_headers",
+ "${tint_spirv_tools_dir}:spvtools_val",
+ ]
+ }
+ }
+ }
}
diff --git a/src/tint/lang/spirv/reader/common/helper_test.h b/src/tint/lang/spirv/reader/common/helper_test.h
new file mode 100644
index 0000000..6a0ae0e
--- /dev/null
+++ b/src/tint/lang/spirv/reader/common/helper_test.h
@@ -0,0 +1,58 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_SPIRV_READER_COMMON_HELPER_TEST_H_
+#define SRC_TINT_LANG_SPIRV_READER_COMMON_HELPER_TEST_H_
+
+#include <string>
+#include <vector>
+
+#include "spirv-tools/libspirv.hpp"
+#include "src/tint/utils/result/result.h"
+
+namespace tint::spirv::reader {
+
+/// 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) {
+ StringStream err;
+ std::vector<uint32_t> binary;
+ spvtools::SpirvTools tools(SPV_ENV_UNIVERSAL_1_0);
+ 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;
+ });
+ if (!tools.Assemble(spirv_asm, &binary, SPV_TEXT_TO_BINARY_OPTION_PRESERVE_NUMERIC_IDS)) {
+ return err.str();
+ }
+ return binary;
+}
+
+} // namespace tint::spirv::reader
+
+#endif // SRC_TINT_LANG_SPIRV_READER_COMMON_HELPER_TEST_H_
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.bazel b/src/tint/lang/spirv/reader/lower/BUILD.bazel
new file mode 100644
index 0000000..01fd572
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/BUILD.bazel
@@ -0,0 +1,106 @@
+# Copyright 2023 The Dawn & Tint Authors
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+#
+# 1. Redistributions of source code must retain the above copyright notice, this
+# list of conditions and the following disclaimer.
+#
+# 2. Redistributions in binary form must reproduce the above copyright notice,
+# this list of conditions and the following disclaimer in the documentation
+# and/or other materials provided with the distribution.
+#
+# 3. Neither the name of the copyright holder nor the names of its
+# contributors may be used to endorse or promote products derived from
+# this software without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+# tools/src/cmd/gen/build/BUILD.bazel.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+# Do not modify this file directly
+################################################################################
+
+load("//src/tint:flags.bzl", "COPTS")
+load("@bazel_skylib//lib:selects.bzl", "selects")
+cc_library(
+ name = "lower",
+ srcs = [
+ "lower.cc",
+ "vector_element_pointer.cc",
+ ],
+ hdrs = [
+ "lower.h",
+ "vector_element_pointer.h",
+ ],
+ deps = [
+ "//src/tint/api/common",
+ "//src/tint/lang/core",
+ "//src/tint/lang/core/constant",
+ "//src/tint/lang/core/intrinsic",
+ "//src/tint/lang/core/ir",
+ "//src/tint/lang/core/type",
+ "//src/tint/utils/containers",
+ "//src/tint/utils/diagnostic",
+ "//src/tint/utils/ice",
+ "//src/tint/utils/id",
+ "//src/tint/utils/macros",
+ "//src/tint/utils/math",
+ "//src/tint/utils/memory",
+ "//src/tint/utils/reflection",
+ "//src/tint/utils/result",
+ "//src/tint/utils/rtti",
+ "//src/tint/utils/symbol",
+ "//src/tint/utils/text",
+ "//src/tint/utils/traits",
+ ],
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
+cc_library(
+ name = "test",
+ alwayslink = True,
+ srcs = [
+ "vector_element_pointer_test.cc",
+ ],
+ deps = [
+ "//src/tint/api/common",
+ "//src/tint/lang/core",
+ "//src/tint/lang/core/constant",
+ "//src/tint/lang/core/intrinsic",
+ "//src/tint/lang/core/ir",
+ "//src/tint/lang/core/ir/transform:test",
+ "//src/tint/lang/core/type",
+ "//src/tint/lang/spirv/reader/lower",
+ "//src/tint/utils/containers",
+ "//src/tint/utils/diagnostic",
+ "//src/tint/utils/ice",
+ "//src/tint/utils/id",
+ "//src/tint/utils/macros",
+ "//src/tint/utils/math",
+ "//src/tint/utils/memory",
+ "//src/tint/utils/reflection",
+ "//src/tint/utils/result",
+ "//src/tint/utils/rtti",
+ "//src/tint/utils/symbol",
+ "//src/tint/utils/text",
+ "//src/tint/utils/traits",
+ "@gtest",
+ ],
+ copts = COPTS,
+ visibility = ["//visibility:public"],
+)
+
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.cmake b/src/tint/lang/spirv/reader/lower/BUILD.cmake
new file mode 100644
index 0000000..5f0e866
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/BUILD.cmake
@@ -0,0 +1,104 @@
+# Copyright 2023 The Dawn & Tint Authors
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+#
+# 1. Redistributions of source code must retain the above copyright notice, this
+# list of conditions and the following disclaimer.
+#
+# 2. Redistributions in binary form must reproduce the above copyright notice,
+# this list of conditions and the following disclaimer in the documentation
+# and/or other materials provided with the distribution.
+#
+# 3. Neither the name of the copyright holder nor the names of its
+# contributors may be used to endorse or promote products derived from
+# this software without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+# tools/src/cmd/gen/build/BUILD.cmake.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+# Do not modify this file directly
+################################################################################
+
+################################################################################
+# Target: tint_lang_spirv_reader_lower
+# Kind: lib
+################################################################################
+tint_add_target(tint_lang_spirv_reader_lower lib
+ lang/spirv/reader/lower/lower.cc
+ lang/spirv/reader/lower/lower.h
+ lang/spirv/reader/lower/vector_element_pointer.cc
+ lang/spirv/reader/lower/vector_element_pointer.h
+)
+
+tint_target_add_dependencies(tint_lang_spirv_reader_lower lib
+ tint_api_common
+ tint_lang_core
+ tint_lang_core_constant
+ tint_lang_core_intrinsic
+ tint_lang_core_ir
+ tint_lang_core_type
+ tint_utils_containers
+ tint_utils_diagnostic
+ tint_utils_ice
+ tint_utils_id
+ tint_utils_macros
+ tint_utils_math
+ tint_utils_memory
+ tint_utils_reflection
+ tint_utils_result
+ tint_utils_rtti
+ tint_utils_symbol
+ tint_utils_text
+ tint_utils_traits
+)
+
+################################################################################
+# Target: tint_lang_spirv_reader_lower_test
+# Kind: test
+################################################################################
+tint_add_target(tint_lang_spirv_reader_lower_test test
+ lang/spirv/reader/lower/vector_element_pointer_test.cc
+)
+
+tint_target_add_dependencies(tint_lang_spirv_reader_lower_test test
+ tint_api_common
+ tint_lang_core
+ tint_lang_core_constant
+ tint_lang_core_intrinsic
+ tint_lang_core_ir
+ tint_lang_core_ir_transform_test
+ tint_lang_core_type
+ tint_lang_spirv_reader_lower
+ tint_utils_containers
+ tint_utils_diagnostic
+ tint_utils_ice
+ tint_utils_id
+ tint_utils_macros
+ tint_utils_math
+ tint_utils_memory
+ tint_utils_reflection
+ tint_utils_result
+ tint_utils_rtti
+ tint_utils_symbol
+ tint_utils_text
+ tint_utils_traits
+)
+
+tint_target_add_external_dependencies(tint_lang_spirv_reader_lower_test test
+ "gtest"
+)
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.gn b/src/tint/lang/spirv/reader/lower/BUILD.gn
new file mode 100644
index 0000000..b5f0342
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/BUILD.gn
@@ -0,0 +1,102 @@
+# Copyright 2023 The Dawn & Tint Authors
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+#
+# 1. Redistributions of source code must retain the above copyright notice, this
+# list of conditions and the following disclaimer.
+#
+# 2. Redistributions in binary form must reproduce the above copyright notice,
+# this list of conditions and the following disclaimer in the documentation
+# and/or other materials provided with the distribution.
+#
+# 3. Neither the name of the copyright holder nor the names of its
+# contributors may be used to endorse or promote products derived from
+# this software without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+# AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+# IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+# DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+# FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+# DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+# SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+# OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+# OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+################################################################################
+# File generated by 'tools/src/cmd/gen' using the template:
+# tools/src/cmd/gen/build/BUILD.gn.tmpl
+#
+# To regenerate run: './tools/run gen'
+#
+# Do not modify this file directly
+################################################################################
+
+import("../../../../../../scripts/tint_overrides_with_defaults.gni")
+
+import("${tint_src_dir}/tint.gni")
+
+if (tint_build_unittests || tint_build_benchmarks) {
+ import("//testing/test.gni")
+}
+
+libtint_source_set("lower") {
+ sources = [
+ "lower.cc",
+ "lower.h",
+ "vector_element_pointer.cc",
+ "vector_element_pointer.h",
+ ]
+ deps = [
+ "${tint_src_dir}/api/common",
+ "${tint_src_dir}/lang/core",
+ "${tint_src_dir}/lang/core/constant",
+ "${tint_src_dir}/lang/core/intrinsic",
+ "${tint_src_dir}/lang/core/ir",
+ "${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/utils/containers",
+ "${tint_src_dir}/utils/diagnostic",
+ "${tint_src_dir}/utils/ice",
+ "${tint_src_dir}/utils/id",
+ "${tint_src_dir}/utils/macros",
+ "${tint_src_dir}/utils/math",
+ "${tint_src_dir}/utils/memory",
+ "${tint_src_dir}/utils/reflection",
+ "${tint_src_dir}/utils/result",
+ "${tint_src_dir}/utils/rtti",
+ "${tint_src_dir}/utils/symbol",
+ "${tint_src_dir}/utils/text",
+ "${tint_src_dir}/utils/traits",
+ ]
+}
+if (tint_build_unittests) {
+ tint_unittests_source_set("unittests") {
+ sources = [ "vector_element_pointer_test.cc" ]
+ deps = [
+ "${tint_src_dir}:gmock_and_gtest",
+ "${tint_src_dir}/api/common",
+ "${tint_src_dir}/lang/core",
+ "${tint_src_dir}/lang/core/constant",
+ "${tint_src_dir}/lang/core/intrinsic",
+ "${tint_src_dir}/lang/core/ir",
+ "${tint_src_dir}/lang/core/ir/transform:unittests",
+ "${tint_src_dir}/lang/core/type",
+ "${tint_src_dir}/lang/spirv/reader/lower",
+ "${tint_src_dir}/utils/containers",
+ "${tint_src_dir}/utils/diagnostic",
+ "${tint_src_dir}/utils/ice",
+ "${tint_src_dir}/utils/id",
+ "${tint_src_dir}/utils/macros",
+ "${tint_src_dir}/utils/math",
+ "${tint_src_dir}/utils/memory",
+ "${tint_src_dir}/utils/reflection",
+ "${tint_src_dir}/utils/result",
+ "${tint_src_dir}/utils/rtti",
+ "${tint_src_dir}/utils/symbol",
+ "${tint_src_dir}/utils/text",
+ "${tint_src_dir}/utils/traits",
+ ]
+ }
+}
diff --git a/src/tint/lang/spirv/reader/lower/lower.cc b/src/tint/lang/spirv/reader/lower/lower.cc
new file mode 100644
index 0000000..25f546c
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/lower.cc
@@ -0,0 +1,54 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/lower/lower.h"
+
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/spirv/reader/lower/vector_element_pointer.h"
+
+namespace tint::spirv::reader {
+
+Result<SuccessType> Lower(core::ir::Module& mod) {
+#define RUN_TRANSFORM(name, ...) \
+ do { \
+ auto result = name(__VA_ARGS__); \
+ if (result != Success) { \
+ return result; \
+ } \
+ } while (false)
+
+ RUN_TRANSFORM(lower::VectorElementPointer, mod);
+
+ if (auto res = core::ir::ValidateAndDumpIfNeeded(mod, "end of lowering from SPIR-V");
+ res != Success) {
+ return res.Failure();
+ }
+
+ return Success;
+}
+
+} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/lower/lower.h b/src/tint/lang/spirv/reader/lower/lower.h
new file mode 100644
index 0000000..0d8e3b2
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/lower.h
@@ -0,0 +1,43 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_SPIRV_READER_LOWER_LOWER_H_
+#define SRC_TINT_LANG_SPIRV_READER_LOWER_LOWER_H_
+
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/utils/result/result.h"
+
+namespace tint::spirv::reader {
+
+/// Lower converts a SPIR-V-dialect IR module to a core-dialect IR module
+/// @param mod the IR module
+/// @return the result of the operation
+Result<SuccessType> Lower(core::ir::Module& mod);
+
+} // namespace tint::spirv::reader
+
+#endif // SRC_TINT_LANG_SPIRV_READER_LOWER_LOWER_H_
diff --git a/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc b/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc
new file mode 100644
index 0000000..ad0d3f3
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc
@@ -0,0 +1,178 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/lower/vector_element_pointer.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+namespace tint::spirv::reader::lower {
+
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+
+/// PIMPL state for the transform.
+struct State {
+ /// The IR module.
+ core::ir::Module& ir;
+
+ /// The IR builder.
+ core::ir::Builder b{ir};
+
+ /// The type manager.
+ core::type::Manager& ty{ir.Types()};
+
+ /// Access is an access instruction and the type of the vector that it produces a pointer to.
+ struct Access {
+ /// The access instruction.
+ core::ir::Access* inst;
+ /// The vector type being accessed.
+ const core::type::Type* type;
+ };
+
+ /// Process the module.
+ void Process() {
+ // Find the access instructions that need to be replaced.
+ Vector<Access, 8> worklist;
+ for (auto* inst : ir.instructions.Objects()) {
+ if (!inst->Alive()) {
+ continue;
+ }
+ if (auto* access = inst->As<core::ir::Access>()) {
+ auto* source_ty = access->Object()->Type();
+ if (!source_ty->Is<core::type::Pointer>()) {
+ continue;
+ }
+ source_ty = source_ty->UnwrapPtr();
+
+ // Step through the indices of the access instruction to check for vector types.
+ for (auto* idx : access->Indices()) {
+ if (source_ty->Is<core::type::Vector>()) {
+ // Found an access that is indexing into a vector pointer.
+ worklist.Push(Access{access, source_ty});
+ break;
+ }
+
+ // Update the current source type based on the next index.
+ if (auto* constant = idx->As<core::ir::Constant>()) {
+ auto i = constant->Value()->ValueAs<u32>();
+ source_ty = source_ty->Element(i);
+ } else {
+ source_ty = source_ty->Elements().type;
+ }
+ }
+ }
+ }
+
+ // Replace the access instructions that we found.
+ for (const auto& access : worklist) {
+ ReplaceAccess(access);
+ }
+ }
+
+ /// Replace an access instruction with {load,store}_vector_element instructions.
+ /// @param access the access instruction to replace
+ void ReplaceAccess(const Access& access) {
+ auto* object = access.inst->Object();
+
+ if (access.inst->Indices().Length() > 1) {
+ // Create a new access instruction that stops at the vector pointer.
+ Vector<core::ir::Value*, 8> partial_indices{access.inst->Indices()};
+ partial_indices.Pop();
+ auto addrspace = object->Type()->As<core::type::Pointer>()->AddressSpace();
+ auto* access_to_vec = b.Access(ty.ptr(addrspace, access.type), object, partial_indices);
+ access_to_vec->InsertBefore(access.inst);
+
+ object = access_to_vec->Result(0);
+ }
+
+ // Replace all uses of the original access instruction.
+ auto* index = access.inst->Indices().Back();
+ ReplaceAccessUses(access.inst, object, index);
+
+ // Destroy the original access instruction.
+ access.inst->Destroy();
+ }
+
+ /// Replace all uses of an access instruction with {load,store}_vector_element instructions.
+ /// @param access the access instruction to replace
+ /// @param object the pointer-to-vector source object
+ /// @param index the index of the vector element
+ void ReplaceAccessUses(core::ir::Access* access,
+ core::ir::Value* object,
+ core::ir::Value* index) {
+ Vector<core::ir::Instruction*, 4> to_destroy;
+ access->Result(0)->ForEachUse([&](core::ir::Usage use) {
+ Switch(
+ use.instruction,
+ [&](core::ir::Load* load) {
+ auto* lve = b.LoadVectorElement(object, index);
+ lve->InsertBefore(load);
+ load->Result(0)->ReplaceAllUsesWith(lve->Result(0));
+ to_destroy.Push(load);
+ },
+ [&](core::ir::Store* store) {
+ auto* sve = b.StoreVectorElement(object, index, store->From());
+ sve->InsertBefore(store);
+ to_destroy.Push(store);
+ },
+ [&](core::ir::Access* noop_access) {
+ TINT_ASSERT(noop_access->Indices().IsEmpty());
+ ReplaceAccessUses(noop_access, object, index);
+ to_destroy.Push(noop_access);
+ },
+ TINT_ICE_ON_NO_MATCH);
+ });
+
+ // Clean up old instructions.
+ for (auto* inst : to_destroy) {
+ inst->Destroy();
+ }
+ }
+};
+
+} // namespace
+
+Result<SuccessType> VectorElementPointer(core::ir::Module& ir) {
+ auto result = ValidateAndDumpIfNeeded(ir, "VectorElementPointer transform",
+ EnumSet<core::ir::Capability>{
+ core::ir::Capability::kAllowVectorElementPointer,
+ });
+ if (result != Success) {
+ return result.Failure();
+ }
+
+ State{ir}.Process();
+
+ return Success;
+}
+
+} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/lower/vector_element_pointer.h b/src/tint/lang/spirv/reader/lower/vector_element_pointer.h
new file mode 100644
index 0000000..e07a30c
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/vector_element_pointer.h
@@ -0,0 +1,48 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#ifndef SRC_TINT_LANG_SPIRV_READER_LOWER_VECTOR_ELEMENT_POINTER_H_
+#define SRC_TINT_LANG_SPIRV_READER_LOWER_VECTOR_ELEMENT_POINTER_H_
+
+#include "src/tint/utils/result/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+}
+
+namespace tint::spirv::reader::lower {
+
+/// VectorElementPointer is a transform that removes pointers to vector elements by replacing access
+/// instructions and their uses.
+/// @param module the module to transform
+/// @returns success or failure
+Result<SuccessType> VectorElementPointer(core::ir::Module& module);
+
+} // namespace tint::spirv::reader::lower
+
+#endif // SRC_TINT_LANG_SPIRV_READER_LOWER_VECTOR_ELEMENT_POINTER_H_
diff --git a/src/tint/lang/spirv/reader/lower/vector_element_pointer_test.cc b/src/tint/lang/spirv/reader/lower/vector_element_pointer_test.cc
new file mode 100644
index 0000000..e7160a5
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/vector_element_pointer_test.cc
@@ -0,0 +1,487 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/lower/vector_element_pointer.h"
+
+#include <utility>
+
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+
+namespace tint::spirv::reader::lower {
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+using SpirvReader_VectorElementPointerTest = core::ir::transform::TransformTest;
+
+TEST_F(SpirvReader_VectorElementPointerTest, NonPointerAccess) {
+ auto* vec = b.FunctionParam("vec", ty.vec4<u32>());
+ auto* foo = b.Function("foo", ty.u32());
+ b.Append(foo->Block(), [&] {
+ auto* access = b.Access<u32>(vec, 2_u);
+ b.Return(foo, access);
+ });
+
+ auto* src = R"(
+%foo = func():u32 -> %b1 {
+ %b1 = block {
+ %2:u32 = access %vec, 2u
+ ret %2
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_VectorElementPointerTest, Access_NoIndices) {
+ auto* foo = b.Function("foo", ty.vec4<u32>());
+ b.Append(foo->Block(), [&] {
+ auto* vec = b.Var<function, vec4<u32>>("vec");
+ auto* access = b.Access<ptr<function, vec4<u32>>>(vec);
+ b.Return(foo, b.Load(access));
+ });
+
+ auto* src = R"(
+%foo = func():vec4<u32> -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ %3:ptr<function, vec4<u32>, read_write> = access %vec
+ %4:vec4<u32> = load %3
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_VectorElementPointerTest, Access_NoIndices_Chain) {
+ auto* foo = b.Function("foo", ty.vec4<u32>());
+ b.Append(foo->Block(), [&] {
+ auto* vec = b.Var<function, vec4<u32>>("vec");
+ auto* access_1 = b.Access<ptr<function, vec4<u32>>>(vec);
+ auto* access_2 = b.Access<ptr<function, vec4<u32>>>(access_1);
+ auto* access_3 = b.Access<ptr<function, vec4<u32>>>(access_2);
+ b.Return(foo, b.Load(access_3));
+ });
+
+ auto* src = R"(
+%foo = func():vec4<u32> -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ %3:ptr<function, vec4<u32>, read_write> = access %vec
+ %4:ptr<function, vec4<u32>, read_write> = access %3
+ %5:ptr<function, vec4<u32>, read_write> = access %4
+ %6:vec4<u32> = load %5
+ ret %6
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_VectorElementPointerTest, Access_Component_NoUse) {
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* vec = b.Var<function, vec4<u32>>("vec");
+ b.Access<ptr<function, u32>>(vec, 2_u);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ %3:ptr<function, u32, read_write> = access %vec, 2u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ ret
+ }
+}
+)";
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_VectorElementPointerTest, Load) {
+ auto* foo = b.Function("foo", ty.u32());
+ b.Append(foo->Block(), [&] {
+ auto* vec = b.Var<function, vec4<u32>>("vec");
+ auto* access = b.Access<ptr<function, u32>>(vec, 2_u);
+ auto* load = b.Load(access);
+ b.Return(foo, load);
+ });
+
+ auto* src = R"(
+%foo = func():u32 -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ %3:ptr<function, u32, read_write> = access %vec, 2u
+ %4:u32 = load %3
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():u32 -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ %3:u32 = load_vector_element %vec, 2u
+ ret %3
+ }
+}
+)";
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_VectorElementPointerTest, Store) {
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* vec = b.Var<function, vec4<u32>>("vec");
+ auto* access = b.Access<ptr<function, u32>>(vec, 2_u);
+ b.Store(access, 42_u);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ %3:ptr<function, u32, read_write> = access %vec, 2u
+ store %3, 42u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ store_vector_element %vec, 2u, 42u
+ ret
+ }
+}
+)";
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_VectorElementPointerTest, AccessBeforeUse) {
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* vec = b.Var<function, vec4<u32>>("vec");
+ auto* access_1 = b.Access<ptr<function, u32>>(vec, 2_u);
+ auto* access_2 = b.Access<ptr<function, u32>>(access_1);
+ b.Store(access_2, 42_u);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ %3:ptr<function, u32, read_write> = access %vec, 2u
+ %4:ptr<function, u32, read_write> = access %3
+ store %4, 42u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ store_vector_element %vec, 2u, 42u
+ ret
+ }
+}
+)";
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_VectorElementPointerTest, MultipleUses) {
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* vec = b.Var<function, vec4<u32>>("vec");
+ auto* access = b.Access<ptr<function, u32>>(vec, 2_u);
+ auto* load = b.Load(access);
+ auto* add = b.Add<u32>(load, 1_u);
+ b.Store(access, add);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ %3:ptr<function, u32, read_write> = access %vec, 2u
+ %4:u32 = load %3
+ %5:u32 = add %4, 1u
+ store %3, %5
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %vec:ptr<function, vec4<u32>, read_write> = var
+ %3:u32 = load_vector_element %vec, 2u
+ %4:u32 = add %3, 1u
+ store_vector_element %vec, 2u, %4
+ ret
+ }
+}
+)";
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_VectorElementPointerTest, ViaMatrix) {
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* mat = b.Var<function, mat4x4<f32>>("mat");
+ auto* access = b.Access<ptr<function, f32>>(mat, 1_u, 2_u);
+ b.Store(access, 42_f);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %mat:ptr<function, mat4x4<f32>, read_write> = var
+ %3:ptr<function, f32, read_write> = access %mat, 1u, 2u
+ store %3, 42.0f
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %mat:ptr<function, mat4x4<f32>, read_write> = var
+ %3:ptr<function, vec4<f32>, read_write> = access %mat, 1u
+ store_vector_element %3, 2u, 42.0f
+ ret
+ }
+}
+)";
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_VectorElementPointerTest, ViaArray) {
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* arr = b.Var<function, array<vec4<f32>, 4>>("arr");
+ auto* access = b.Access<ptr<function, f32>>(arr, 1_u, 2_u);
+ b.Store(access, 42_f);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %arr:ptr<function, array<vec4<f32>, 4>, read_write> = var
+ %3:ptr<function, f32, read_write> = access %arr, 1u, 2u
+ store %3, 42.0f
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %arr:ptr<function, array<vec4<f32>, 4>, read_write> = var
+ %3:ptr<function, vec4<f32>, read_write> = access %arr, 1u
+ store_vector_element %3, 2u, 42.0f
+ ret
+ }
+}
+)";
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_VectorElementPointerTest, ViaStruct) {
+ auto* str_ty = ty.Struct(mod.symbols.New("str"), {{
+ mod.symbols.New("vec"),
+ ty.vec4<f32>(),
+ }});
+
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* str = b.Var("str", ty.ptr<function>(str_ty));
+ auto* access = b.Access<ptr<function, f32>>(str, 0_u, 2_u);
+ b.Store(access, 42_f);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+str = struct @align(16) {
+ vec:vec4<f32> @offset(0)
+}
+
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %str:ptr<function, str, read_write> = var
+ %3:ptr<function, f32, read_write> = access %str, 0u, 2u
+ store %3, 42.0f
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+str = struct @align(16) {
+ vec:vec4<f32> @offset(0)
+}
+
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %str:ptr<function, str, read_write> = var
+ %3:ptr<function, vec4<f32>, read_write> = access %str, 0u
+ store_vector_element %3, 2u, 42.0f
+ ret
+ }
+}
+)";
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_VectorElementPointerTest, DeeplyNested) {
+ auto* inner_arr = ty.array(ty.mat4x4<f32>(), 4);
+ auto* str_ty = ty.Struct(mod.symbols.New("str"), {{
+ mod.symbols.New("inner"),
+ inner_arr,
+ }});
+ auto* outer_arr = ty.array(str_ty, 4);
+
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* arr = b.Var("arr", ty.ptr<function>(outer_arr));
+ auto* access = b.Access<ptr<function, f32>>(arr, 1_u, 0_u, 3_u, 2_u, 1_u);
+ b.Store(access, 42_f);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+str = struct @align(16) {
+ inner:array<mat4x4<f32>, 4> @offset(0)
+}
+
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %arr:ptr<function, array<str, 4>, read_write> = var
+ %3:ptr<function, f32, read_write> = access %arr, 1u, 0u, 3u, 2u, 1u
+ store %3, 42.0f
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+str = struct @align(16) {
+ inner:array<mat4x4<f32>, 4> @offset(0)
+}
+
+%foo = func():void -> %b1 {
+ %b1 = block {
+ %arr:ptr<function, array<str, 4>, read_write> = var
+ %3:ptr<function, vec4<f32>, read_write> = access %arr, 1u, 0u, 3u, 2u
+ store_vector_element %3, 1u, 42.0f
+ ret
+ }
+}
+)";
+
+ Run(VectorElementPointer);
+
+ EXPECT_EQ(expect, str());
+}
+
+} // namespace
+} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.bazel b/src/tint/lang/spirv/reader/parser/BUILD.bazel
index d41f46f..2bfaad6 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/parser/BUILD.bazel
@@ -108,6 +108,7 @@
"@gtest",
] + select({
":tint_build_spv_reader": [
+ "//src/tint/lang/spirv/reader/common:test",
"//src/tint/lang/spirv/reader/parser",
],
"//conditions:default": [],
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.cmake b/src/tint/lang/spirv/reader/parser/BUILD.cmake
index e82f651..c8abe87 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/parser/BUILD.cmake
@@ -120,6 +120,7 @@
if(TINT_BUILD_SPV_READER)
tint_target_add_dependencies(tint_lang_spirv_reader_parser_test test
+ tint_lang_spirv_reader_common_test
tint_lang_spirv_reader_parser
)
endif(TINT_BUILD_SPV_READER)
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.gn b/src/tint/lang/spirv/reader/parser/BUILD.gn
index 430c576..26c2954 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/parser/BUILD.gn
@@ -116,7 +116,10 @@
]
if (tint_build_spv_reader) {
- deps += [ "${tint_src_dir}/lang/spirv/reader/parser" ]
+ deps += [
+ "${tint_src_dir}/lang/spirv/reader/common:unittests",
+ "${tint_src_dir}/lang/spirv/reader/parser",
+ ]
}
if (tint_build_spv_reader || tint_build_spv_writer) {
diff --git a/src/tint/lang/spirv/reader/parser/helper_test.h b/src/tint/lang/spirv/reader/parser/helper_test.h
index decab00..c7fc02a 100644
--- a/src/tint/lang/spirv/reader/parser/helper_test.h
+++ b/src/tint/lang/spirv/reader/parser/helper_test.h
@@ -34,10 +34,10 @@
#include "gmock/gmock.h"
#include "gtest/gtest.h"
-#include "spirv-tools/libspirv.hpp"
#include "src/tint/lang/core/ir/disassembler.h"
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/spirv/reader/common/helper_test.h"
#include "src/tint/lang/spirv/reader/parser/parser.h"
namespace tint::spirv::reader {
@@ -56,23 +56,6 @@
template <typename BASE>
class SpirvParserTestHelperBase : public BASE {
protected:
- /// 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
- static Result<std::vector<uint32_t>, std::string> Assemble(std::string spirv_asm) {
- StringStream err;
- std::vector<uint32_t> binary;
- spvtools::SpirvTools tools(SPV_ENV_UNIVERSAL_1_0);
- 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;
- });
- if (!tools.Assemble(spirv_asm, &binary, SPV_TEXT_TO_BINARY_OPTION_PRESERVE_NUMERIC_IDS)) {
- return err.str();
- }
- return binary;
- }
-
/// 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
@@ -89,8 +72,11 @@
return parsed.Failure();
}
- // Validate the IR module.
- auto validated = core::ir::Validate(parsed.Get());
+ // Validate the IR module against the capabilities supported by the SPIR-V dialect.
+ auto validated =
+ core::ir::Validate(parsed.Get(), EnumSet<core::ir::Capability>{
+ core::ir::Capability::kAllowVectorElementPointer,
+ });
if (validated != Success) {
return validated.Failure();
}
diff --git a/src/tint/lang/spirv/reader/parser/memory_test.cc b/src/tint/lang/spirv/reader/parser/memory_test.cc
index 2031696..9f8c75f 100644
--- a/src/tint/lang/spirv/reader/parser/memory_test.cc
+++ b/src/tint/lang/spirv/reader/parser/memory_test.cc
@@ -83,8 +83,7 @@
)");
}
-// TODO(jrprice): We need to handle pointer-to-vector component somewhere.
-TEST_F(SpirvParserTest, DISABLED_Load_VectorComponent) {
+TEST_F(SpirvParserTest, Load_VectorComponent) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -350,8 +349,7 @@
)");
}
-// TODO(jrprice): We need to handle pointer-to-vector component somewhere.
-TEST_F(SpirvParserTest, DISABLED_Store_VectorComponent) {
+TEST_F(SpirvParserTest, Store_VectorComponent) {
EXPECT_IR(R"(
OpCapability Shader
OpMemoryModel Logical GLSL450
@@ -359,7 +357,7 @@
OpExecutionMode %main LocalSize 1 1 1
%void = OpTypeVoid
%u32 = OpTypeInt 32 0
- %u32_42 = OpConstantInt %u32 42
+ %u32_42 = OpConstant %u32 42
%vec4u = OpTypeVector %u32 4
%u32_ptr = OpTypePointer Function %u32
%vec4u_ptr = OpTypePointer Function %vec4u
diff --git a/src/tint/lang/spirv/reader/reader.cc b/src/tint/lang/spirv/reader/reader.cc
index 11199ea..d70c16c 100644
--- a/src/tint/lang/spirv/reader/reader.cc
+++ b/src/tint/lang/spirv/reader/reader.cc
@@ -31,17 +31,22 @@
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/spirv/reader/ast_parser/parse.h"
+#include "src/tint/lang/spirv/reader/lower/lower.h"
#include "src/tint/lang/spirv/reader/parser/parser.h"
namespace tint::spirv::reader {
Result<core::ir::Module> ReadIR(const std::vector<uint32_t>& input) {
+ // Parse the input SPIR-V to the SPIR-V dialect of the IR.
auto mod = Parse(Slice(input.data(), input.size()));
if (mod != Success) {
return mod.Failure();
}
- // TODO(crbug.com/tint/1907): Lower the module to core dialect.
+ // Lower the module to the core dialect of the IR.
+ if (auto res = Lower(mod.Get()); res != Success) {
+ return std::move(res.Failure());
+ }
return mod;
}
diff --git a/src/tint/lang/spirv/reader/reader_test.cc b/src/tint/lang/spirv/reader/reader_test.cc
new file mode 100644
index 0000000..1a086a8
--- /dev/null
+++ b/src/tint/lang/spirv/reader/reader_test.cc
@@ -0,0 +1,139 @@
+// Copyright 2024 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/reader.h"
+
+#include <string>
+
+#include "gmock/gmock.h"
+#include "gtest/gtest.h"
+#include "src/tint/lang/core/ir/disassembler.h"
+#include "src/tint/lang/core/ir/module.h"
+#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/spirv/reader/common/helper_test.h"
+
+namespace tint::spirv::reader {
+namespace {
+
+class SpirvReaderTest : public testing::Test {
+ protected:
+ /// Run the reader on a SPIR-V module and return the Tint IR or an error string.
+ /// @param spirv_asm the SPIR-V assembly to read
+ /// @returns the disassembled Tint IR or an error
+ Result<std::string> Run(std::string spirv_asm) {
+ // Assemble the SPIR-V input.
+ auto binary = Assemble(spirv_asm);
+ if (binary != Success) {
+ return binary.Failure();
+ }
+
+ // Read the SPIR-V to produce a core IR module.
+ auto ir = ReadIR(binary.Get());
+ if (ir != Success) {
+ return ir.Failure();
+ }
+
+ // Validate the IR module against the core dialect.
+ auto validated = core::ir::Validate(ir.Get());
+ if (validated != Success) {
+ return validated.Failure();
+ }
+
+ // Return the disassembled IR module.
+ return "\n" + core::ir::Disassemble(ir.Get());
+ }
+};
+
+TEST_F(SpirvReaderTest, Load_VectorComponent) {
+ auto got = Run(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %u32 = OpTypeInt 32 0
+ %vec4u = OpTypeVector %u32 4
+ %u32_ptr = OpTypePointer Function %u32
+ %vec4u_ptr = OpTypePointer Function %vec4u
+ %u32_2 = OpConstant %u32 2
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %var = OpVariable %vec4u_ptr Function
+ %access = OpAccessChain %u32_ptr %var %u32_2
+ %load = OpLoad %u32 %access
+ OpReturn
+ OpFunctionEnd
+)");
+ ASSERT_EQ(got, Success);
+ EXPECT_EQ(got, R"(
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %2:ptr<function, vec4<u32>, read_write> = var
+ %3:u32 = load_vector_element %2, 2u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvReaderTest, Store_VectorComponent) {
+ auto got = Run(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %u32 = OpTypeInt 32 0
+ %u32_42 = OpConstant %u32 42
+ %vec4u = OpTypeVector %u32 4
+ %u32_ptr = OpTypePointer Function %u32
+ %vec4u_ptr = OpTypePointer Function %vec4u
+ %u32_2 = OpConstant %u32 2
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ %var = OpVariable %vec4u_ptr Function
+ %access = OpAccessChain %u32_ptr %var %u32_2
+ OpStore %access %u32_42
+ OpReturn
+ OpFunctionEnd
+)");
+ ASSERT_EQ(got, Success);
+ EXPECT_EQ(got, R"(
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
+ %b1 = block {
+ %2:ptr<function, vec4<u32>, read_write> = var
+ store_vector_element %2, 2u, 42u
+ ret
+ }
+}
+)");
+}
+
+} // namespace
+} // namespace tint::spirv::reader
