[spirv-reader] Add transform to handle shader IO
Move module-scope output variables to the private address space, wrap
entry points that use them, and then copy their values to the return
value of the wrapper function.
Replace module-scope input variables with function parameters and pass
them down the call stack from the entry point.
Modify the type of sample_mask builtins to a scalar u32.
Bug: 42250952
Change-Id: I01c958803ec4f88c1f2dbb21fdf2206501145176
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/169223
Reviewed-by: Ben Clayton <bclayton@google.com>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.bazel b/src/tint/lang/spirv/reader/lower/BUILD.bazel
index 01fd572..1348c8d 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/lower/BUILD.bazel
@@ -40,10 +40,12 @@
name = "lower",
srcs = [
"lower.cc",
+ "shader_io.cc",
"vector_element_pointer.cc",
],
hdrs = [
"lower.h",
+ "shader_io.h",
"vector_element_pointer.h",
],
deps = [
@@ -52,6 +54,7 @@
"//src/tint/lang/core/constant",
"//src/tint/lang/core/intrinsic",
"//src/tint/lang/core/ir",
+ "//src/tint/lang/core/ir/transform/common",
"//src/tint/lang/core/type",
"//src/tint/utils/containers",
"//src/tint/utils/diagnostic",
@@ -74,6 +77,7 @@
name = "test",
alwayslink = True,
srcs = [
+ "shader_io_test.cc",
"vector_element_pointer_test.cc",
],
deps = [
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.cmake b/src/tint/lang/spirv/reader/lower/BUILD.cmake
index 5f0e866..98de660 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/lower/BUILD.cmake
@@ -41,6 +41,8 @@
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/shader_io.cc
+ lang/spirv/reader/lower/shader_io.h
lang/spirv/reader/lower/vector_element_pointer.cc
lang/spirv/reader/lower/vector_element_pointer.h
)
@@ -51,6 +53,7 @@
tint_lang_core_constant
tint_lang_core_intrinsic
tint_lang_core_ir
+ tint_lang_core_ir_transform_common
tint_lang_core_type
tint_utils_containers
tint_utils_diagnostic
@@ -72,6 +75,7 @@
# Kind: test
################################################################################
tint_add_target(tint_lang_spirv_reader_lower_test test
+ lang/spirv/reader/lower/shader_io_test.cc
lang/spirv/reader/lower/vector_element_pointer_test.cc
)
diff --git a/src/tint/lang/spirv/reader/lower/BUILD.gn b/src/tint/lang/spirv/reader/lower/BUILD.gn
index b5f0342..37c4709 100644
--- a/src/tint/lang/spirv/reader/lower/BUILD.gn
+++ b/src/tint/lang/spirv/reader/lower/BUILD.gn
@@ -46,6 +46,8 @@
sources = [
"lower.cc",
"lower.h",
+ "shader_io.cc",
+ "shader_io.h",
"vector_element_pointer.cc",
"vector_element_pointer.h",
]
@@ -55,6 +57,7 @@
"${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/common",
"${tint_src_dir}/lang/core/type",
"${tint_src_dir}/utils/containers",
"${tint_src_dir}/utils/diagnostic",
@@ -73,7 +76,10 @@
}
if (tint_build_unittests) {
tint_unittests_source_set("unittests") {
- sources = [ "vector_element_pointer_test.cc" ]
+ sources = [
+ "shader_io_test.cc",
+ "vector_element_pointer_test.cc",
+ ]
deps = [
"${tint_src_dir}:gmock_and_gtest",
"${tint_src_dir}/api/common",
diff --git a/src/tint/lang/spirv/reader/lower/lower.cc b/src/tint/lang/spirv/reader/lower/lower.cc
index 25f546c..433d0cd 100644
--- a/src/tint/lang/spirv/reader/lower/lower.cc
+++ b/src/tint/lang/spirv/reader/lower/lower.cc
@@ -28,6 +28,7 @@
#include "src/tint/lang/spirv/reader/lower/lower.h"
#include "src/tint/lang/core/ir/validator.h"
+#include "src/tint/lang/spirv/reader/lower/shader_io.h"
#include "src/tint/lang/spirv/reader/lower/vector_element_pointer.h"
namespace tint::spirv::reader {
@@ -42,6 +43,7 @@
} while (false)
RUN_TRANSFORM(lower::VectorElementPointer, mod);
+ RUN_TRANSFORM(lower::ShaderIO, mod);
if (auto res = core::ir::ValidateAndDumpIfNeeded(mod, "end of lowering from SPIR-V");
res != Success) {
diff --git a/src/tint/lang/spirv/reader/lower/shader_io.cc b/src/tint/lang/spirv/reader/lower/shader_io.cc
new file mode 100644
index 0000000..a786791
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/shader_io.cc
@@ -0,0 +1,426 @@
+// 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/shader_io.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/transform/common/referenced_module_vars.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()};
+
+ /// A map from block to its containing function.
+ Hashmap<core::ir::Block*, core::ir::Function*, 64> block_to_function{};
+
+ /// A map from each function to a map from input variable to parameter.
+ Hashmap<core::ir::Function*, Hashmap<core::ir::Var*, core::ir::Value*, 4>, 8>
+ function_parameter_map{};
+
+ /// The set of output variables that have been processed.
+ Hashset<core::ir::Var*, 4> output_variables{};
+
+ /// The mapping from functions to their transitively referenced output variables.
+ core::ir::ReferencedModuleVars referenced_output_vars{
+ ir, [](const core::ir::Var* var) {
+ auto* view = var->Result(0)->Type()->As<core::type::MemoryView>();
+ return view && view->AddressSpace() == core::AddressSpace::kOut;
+ }};
+
+ /// Process the module.
+ void Process() {
+ // Process outputs first, as that may introduce new functions that input variables need to
+ // be propagated through.
+ ProcessOutputs();
+ ProcessInputs();
+ }
+
+ /// Process output variables.
+ /// Changes output variables to the `private` address space and wraps entry points that produce
+ /// outputs with new functions that copy the outputs from the private variables to the return
+ /// value.
+ void ProcessOutputs() {
+ // Update entry point functions to return their outputs, using a wrapper function.
+ // Use a worklist as `ProcessEntryPointOutputs()` will add new functions.
+ Vector<core::ir::Function*, 4> entry_points;
+ for (auto& func : ir.functions) {
+ if (func->Stage() != core::ir::Function::PipelineStage::kUndefined) {
+ entry_points.Push(func);
+ }
+ }
+ for (auto& ep : entry_points) {
+ ProcessEntryPointOutputs(ep);
+ }
+
+ // Remove attributes from all of the original structs and module-scope output variables.
+ // This is done last as we need to copy attributes during `ProcessEntryPointOutputs()`.
+ for (auto& var : output_variables) {
+ var->SetAttributes({});
+ if (auto* str = var->Result(0)->Type()->UnwrapPtr()->As<core::type::Struct>()) {
+ for (auto* member : str->Members()) {
+ // TODO(crbug.com/tint/745): Remove the const_cast.
+ const_cast<core::type::StructMember*>(member)->SetAttributes({});
+ }
+ }
+ }
+ }
+
+ /// Process input variables.
+ /// Pass inputs down the call stack as parameters to any functions that need them.
+ void ProcessInputs() {
+ // Seed the block-to-function map with the function entry blocks.
+ for (auto& func : ir.functions) {
+ block_to_function.Add(func->Block(), func);
+ }
+
+ // Gather the list of all module-scope input variables.
+ Vector<core::ir::Var*, 4> inputs;
+ for (auto* global : *ir.root_block) {
+ if (auto* var = global->As<core::ir::Var>()) {
+ auto addrspace = var->Result(0)->Type()->As<core::type::Pointer>()->AddressSpace();
+ if (addrspace == core::AddressSpace::kIn) {
+ inputs.Push(var);
+ }
+ }
+ }
+
+ // Replace the input variables with function parameters.
+ for (auto* var : inputs) {
+ ReplaceInputPointerUses(var, var->Result(0));
+ var->Destroy();
+ }
+ }
+
+ /// Replace an output pointer address space to make it `private`.
+ /// @param value the output variable
+ void ReplaceOutputPointerAddressSpace(core::ir::InstructionResult* value) {
+ // Change the address space to `private`.
+ auto* old_ptr_type = value->Type();
+ auto* new_ptr_type = ty.ptr(core::AddressSpace::kPrivate, old_ptr_type->UnwrapPtr());
+ value->SetType(new_ptr_type);
+
+ // Update all uses of the module-scope variable.
+ value->ForEachUse([&](core::ir::Usage use) {
+ if (auto* access = use.instruction->As<core::ir::Access>()) {
+ ReplaceOutputPointerAddressSpace(access->Result(0));
+ } else if (!use.instruction->IsAnyOf<core::ir::Load, core::ir::LoadVectorElement,
+ core::ir::Store, core::ir::StoreVectorElement>()) {
+ TINT_UNREACHABLE()
+ << "unexpected instruction: " << use.instruction->TypeInfo().name;
+ }
+ });
+ }
+
+ /// Process the outputs of an entry point function, adding a wrapper function to forward outputs
+ /// through the return value.
+ /// @param ep the entry point
+ void ProcessEntryPointOutputs(core::ir::Function* ep) {
+ const auto& referenced_outputs = referenced_output_vars.TransitiveReferences(ep);
+ if (referenced_outputs.IsEmpty()) {
+ return;
+ }
+
+ // Add a wrapper function to return either a single value or a struct.
+ auto* wrapper = b.Function(ty.void_(), ep->Stage());
+ if (auto name = ir.NameOf(ep)) {
+ ir.SetName(ep, name.Name() + "_inner");
+ ir.SetName(wrapper, name);
+ }
+
+ // Call the original entry point and make it a regular function.
+ ep->SetStage(core::ir::Function::PipelineStage::kUndefined);
+ b.Append(wrapper->Block(), [&] { //
+ b.Call(ep);
+ });
+
+ // Collect all outputs into a list of struct member declarations.
+ // Also add instructions to load their final values in the wrapper function.
+ Vector<core::ir::Value*, 4> results;
+ Vector<core::type::Manager::StructMemberDesc, 4> output_descriptors;
+ auto add_output = [&](Symbol name, const core::type::Type* type,
+ core::type::StructMemberAttributes attributes) {
+ if (!name) {
+ name = ir.symbols.New();
+ }
+ output_descriptors.Push(core::type::Manager::StructMemberDesc{name, type, attributes});
+ };
+ for (auto* var : referenced_outputs) {
+ // Change the address space of the variable to private and update its uses, if we
+ // haven't already seen this variable.
+ if (output_variables.Add(var)) {
+ ReplaceOutputPointerAddressSpace(var->Result(0));
+ }
+
+ // Copy the variable attributes to the struct member.
+ const auto& original_attributes = var->Attributes();
+ core::type::StructMemberAttributes var_attributes;
+ var_attributes.invariant = original_attributes.invariant;
+ var_attributes.builtin = original_attributes.builtin;
+ var_attributes.location = original_attributes.location;
+ var_attributes.interpolation = original_attributes.interpolation;
+
+ auto var_type = var->Result(0)->Type()->UnwrapPtr();
+ if (auto* str = var_type->As<core::type::Struct>()) {
+ // Add an output for each member of the struct.
+ for (auto* member : str->Members()) {
+ // Use the base variable attributes if not specified directly on the member.
+ auto member_attributes = member->Attributes();
+ if (auto base_loc = var_attributes.location) {
+ // Location values increment from the base location value on the variable.
+ member_attributes.location = base_loc.value() + member->Index();
+ }
+ if (!member_attributes.interpolation) {
+ member_attributes.interpolation = var_attributes.interpolation;
+ }
+
+ add_output(member->Name(), member->Type(), std::move(member_attributes));
+
+ // Load the final result from the member of the original struct variable.
+ b.Append(wrapper->Block(), [&] { //
+ auto* access =
+ b.Access(ty.ptr<private_>(member->Type()), var, u32(member->Index()));
+ results.Push(b.Load(access)->Result(0));
+ });
+ }
+ } else {
+ // Load the final result from the original variable.
+ b.Append(wrapper->Block(), [&] {
+ results.Push(b.Load(var)->Result(0));
+
+ // If we're dealing with sample_mask, extract the scalar from the array.
+ if (var_attributes.builtin == core::BuiltinValue::kSampleMask) {
+ var_type = ty.u32();
+ results.Back() = b.Access(ty.u32(), results.Back(), u32(0))->Result(0);
+ }
+ });
+ add_output(ir.NameOf(var), var_type, std::move(var_attributes));
+ }
+ }
+
+ if (output_descriptors.Length() == 1) {
+ // Copy the output attributes to the function return.
+ const auto& attributes = output_descriptors[0].attributes;
+ wrapper->SetReturnInvariant(attributes.invariant);
+ if (attributes.builtin) {
+ wrapper->SetReturnBuiltin(attributes.builtin.value());
+ } else if (attributes.location) {
+ core::ir::Location loc;
+ loc.value = attributes.location.value();
+ loc.interpolation = attributes.interpolation;
+ wrapper->SetReturnLocation(std::move(loc));
+ }
+
+ // Return the output from the wrapper function.
+ wrapper->SetReturnType(output_descriptors[0].type);
+ b.Append(wrapper->Block(), [&] { //
+ b.Return(wrapper, results[0]);
+ });
+ } else {
+ // Create a struct to hold all of the output values.
+ auto* str = ty.Struct(ir.symbols.New(), std::move(output_descriptors));
+ wrapper->SetReturnType(str);
+
+ // Collect the output values and return them from the wrapper function.
+ b.Append(wrapper->Block(), [&] { //
+ b.Return(wrapper, b.Construct(str, std::move(results)));
+ });
+ }
+ }
+
+ /// Replace a use of an input pointer value.
+ /// @param var the originating input variable
+ /// @param value the input pointer value
+ void ReplaceInputPointerUses(core::ir::Var* var, core::ir::Value* value) {
+ Vector<core::ir::Instruction*, 8> to_destroy;
+ value->ForEachUse([&](core::ir::Usage use) {
+ auto* object = value;
+ if (object->Type()->Is<core::type::Pointer>()) {
+ // Get (or create) the function parameter that will replace the variable.
+ auto* func = ContainingFunction(use.instruction);
+ object = GetParameter(func, var);
+ }
+
+ Switch(
+ use.instruction,
+ [&](core::ir::Load* l) {
+ // Fold the load away and replace its uses with the new parameter.
+ l->Result(0)->ReplaceAllUsesWith(object);
+ to_destroy.Push(l);
+ },
+ [&](core::ir::LoadVectorElement* lve) {
+ // Replace the vector element load with an access instruction.
+ auto* access = b.AccessWithResult(lve->DetachResult(), object, lve->Index());
+ access->InsertBefore(lve);
+ to_destroy.Push(lve);
+ },
+ [&](core::ir::Access* a) {
+ if (!a->Indices().IsEmpty()) {
+ // Remove the pointer from the source and destination type.
+ a->SetOperand(core::ir::Access::kObjectOperandOffset, object);
+ a->Result(0)->SetType(a->Result(0)->Type()->UnwrapPtr());
+ ReplaceInputPointerUses(var, a->Result(0));
+ } else {
+ // Fold the access away and replace its uses.
+ ReplaceInputPointerUses(var, a->Result(0));
+ to_destroy.Push(a);
+ }
+ },
+ TINT_ICE_ON_NO_MATCH);
+ });
+
+ // Clean up orphaned instructions.
+ for (auto* inst : to_destroy) {
+ inst->Destroy();
+ }
+ }
+
+ /// Get the function that contains an instruction.
+ /// @param inst the instruction
+ /// @returns the function
+ core::ir::Function* ContainingFunction(core::ir::Instruction* inst) {
+ return block_to_function.GetOrAdd(inst->Block(), [&] { //
+ return ContainingFunction(inst->Block()->Parent());
+ });
+ }
+
+ /// Get or create a function parameter to replace a module-scope variable.
+ /// @param func the function
+ /// @param var the module-scope variable
+ /// @returns the function parameter
+ core::ir::Value* GetParameter(core::ir::Function* func, core::ir::Var* var) {
+ return function_parameter_map.GetOrAddZero(func).GetOrAdd(var, [&] {
+ const bool entry_point = func->Stage() != core::ir::Function::PipelineStage::kUndefined;
+ auto* var_type = var->Result(0)->Type()->UnwrapPtr();
+
+ // Use a scalar u32 for sample_mask builtins for entry point parameters.
+ if (entry_point && var->Attributes().builtin == core::BuiltinValue::kSampleMask) {
+ TINT_ASSERT(var_type->Is<core::type::Array>());
+ TINT_ASSERT(var_type->As<core::type::Array>()->ConstantCount() == 1u);
+ var_type = ty.u32();
+ }
+
+ // Create a new function parameter for the input.
+ auto* param = b.FunctionParam(var_type);
+ func->AppendParam(param);
+ if (auto name = ir.NameOf(var)) {
+ ir.SetName(param, name);
+ }
+
+ // Add attributes to the parameter if this is an entry point function.
+ if (entry_point) {
+ AddEntryPointParameterAttributes(param, var->Attributes());
+ }
+
+ // Update the callsites of this function.
+ func->ForEachUse([&](core::ir::Usage use) {
+ if (auto* call = use.instruction->As<core::ir::UserCall>()) {
+ // Recurse into the calling function.
+ auto* caller = ContainingFunction(call);
+ call->AppendArg(GetParameter(caller, var));
+ } else if (!use.instruction->Is<core::ir::Return>()) {
+ TINT_UNREACHABLE()
+ << "unexpected instruction: " << use.instruction->TypeInfo().name;
+ }
+ });
+
+ core::ir::Value* result = param;
+ if (entry_point && var->Attributes().builtin == core::BuiltinValue::kSampleMask) {
+ // Construct an array from the scalar sample_mask builtin value for entry points.
+ b.Prepend(func->Block(), [&] { //
+ result = b.Construct(var->Result(0)->Type()->UnwrapPtr(), param)->Result(0);
+ });
+ }
+ return result;
+ });
+ }
+
+ /// Add attributes to an entry point function parameter.
+ /// @param param the parameter
+ /// @param attributes the attributes
+ void AddEntryPointParameterAttributes(core::ir::FunctionParam* param,
+ const core::ir::IOAttributes& attributes) {
+ if (auto* str = param->Type()->UnwrapPtr()->As<core::type::Struct>()) {
+ for (auto* member : str->Members()) {
+ // Use the base variable attributes if not specified directly on the member.
+ auto member_attributes = member->Attributes();
+ if (auto base_loc = attributes.location) {
+ // Location values increment from the base location value on the variable.
+ member_attributes.location = base_loc.value() + member->Index();
+ }
+ if (!member_attributes.interpolation) {
+ member_attributes.interpolation = attributes.interpolation;
+ }
+ // TODO(crbug.com/tint/745): Remove the const_cast.
+ const_cast<core::type::StructMember*>(member)->SetAttributes(
+ std::move(member_attributes));
+ }
+ } else {
+ // Set attributes directly on the function parameter.
+ param->SetInvariant(attributes.invariant);
+ if (attributes.builtin) {
+ param->SetBuiltin(attributes.builtin.value());
+ } else if (attributes.location) {
+ core::ir::Location loc;
+ loc.value = attributes.location.value();
+ loc.interpolation = attributes.interpolation;
+ param->SetLocation(std::move(loc));
+ }
+ }
+ }
+};
+
+} // namespace
+
+Result<SuccessType> ShaderIO(core::ir::Module& ir) {
+ auto result = ValidateAndDumpIfNeeded(ir, "ShaderIO transform");
+ 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/shader_io.h b/src/tint/lang/spirv/reader/lower/shader_io.h
new file mode 100644
index 0000000..8347585
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/shader_io.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_SHADER_IO_H_
+#define SRC_TINT_LANG_SPIRV_READER_LOWER_SHADER_IO_H_
+
+#include "src/tint/utils/result/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+}
+
+namespace tint::spirv::reader::lower {
+
+/// ShaderIO is a transform that converts SPIR-V's style of shader IO (using global variables) into
+/// the form expected by Tint's core IR (using function parameters and return values).
+/// @param module the module to transform
+/// @returns success or failure
+Result<SuccessType> ShaderIO(core::ir::Module& module);
+
+} // namespace tint::spirv::reader::lower
+
+#endif // SRC_TINT_LANG_SPIRV_READER_LOWER_SHADER_IO_H_
diff --git a/src/tint/lang/spirv/reader/lower/shader_io_test.cc b/src/tint/lang/spirv/reader/lower/shader_io_test.cc
new file mode 100644
index 0000000..d004762
--- /dev/null
+++ b/src/tint/lang/spirv/reader/lower/shader_io_test.cc
@@ -0,0 +1,2267 @@
+// 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/shader_io.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
+
+class SpirvReader_ShaderIOTest : public core::ir::transform::TransformTest {
+ protected:
+ core::type::StructMemberAttributes BuiltinAttrs(core::BuiltinValue builtin) {
+ core::type::StructMemberAttributes attrs;
+ attrs.builtin = builtin;
+ return attrs;
+ }
+ core::type::StructMemberAttributes LocationAttrs(
+ uint32_t location,
+ std::optional<core::Interpolation> interpolation = std::nullopt) {
+ core::type::StructMemberAttributes attrs;
+ attrs.location = location;
+ attrs.interpolation = interpolation;
+ return attrs;
+ }
+};
+
+TEST_F(SpirvReader_ShaderIOTest, NoInputsOrOutputs) {
+ auto* ep = b.Function("foo", ty.void_());
+ ep->SetStage(core::ir::Function::PipelineStage::kCompute);
+
+ b.Append(ep->Block(), [&] { //
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute func():void {
+ $B1: {
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Inputs) {
+ auto* front_facing = b.Var("front_facing", ty.ptr(core::AddressSpace::kIn, ty.bool_()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kFrontFacing;
+ front_facing->SetAttributes(std::move(attributes));
+ }
+ auto* position = b.Var("position", ty.ptr(core::AddressSpace::kIn, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kPosition;
+ attributes.invariant = true;
+ position->SetAttributes(std::move(attributes));
+ }
+ auto* color1 = b.Var("color1", ty.ptr(core::AddressSpace::kIn, ty.f32()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 0;
+ color1->SetAttributes(std::move(attributes));
+ }
+ auto* color2 = b.Var("color2", ty.ptr(core::AddressSpace::kIn, ty.f32()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1;
+ attributes.interpolation = core::Interpolation{core::InterpolationType::kLinear,
+ core::InterpolationSampling::kSample};
+ color2->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(front_facing);
+ mod.root_block->Append(position);
+ mod.root_block->Append(color1);
+ mod.root_block->Append(color2);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] {
+ auto* ifelse = b.If(b.Load(front_facing));
+ b.Append(ifelse->True(), [&] {
+ auto* position_value = b.Load(position);
+ auto* color1_value = b.Load(color1);
+ auto* color2_value = b.Load(color2);
+ b.Multiply(ty.vec4<f32>(), position_value, b.Add(ty.f32(), color1_value, color2_value));
+ b.ExitIf(ifelse);
+ });
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %front_facing:ptr<__in, bool, read> = var @builtin(front_facing)
+ %position:ptr<__in, vec4<f32>, read> = var @invariant @builtin(position)
+ %color1:ptr<__in, f32, read> = var @location(0)
+ %color2:ptr<__in, f32, read> = var @location(1) @interpolate(linear, sample)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ %6:bool = load %front_facing
+ if %6 [t: $B3] { # if_1
+ $B3: { # true
+ %7:vec4<f32> = load %position
+ %8:f32 = load %color1
+ %9:f32 = load %color2
+ %10:f32 = add %8, %9
+ %11:vec4<f32> = mul %7, %10
+ exit_if # if_1
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = @fragment func(%front_facing:bool [@front_facing], %position:vec4<f32> [@invariant, @position], %color1:f32 [@location(0)], %color2:f32 [@location(1), @interpolate(linear, sample)]):void {
+ $B1: {
+ if %front_facing [t: $B2] { # if_1
+ $B2: { # true
+ %6:f32 = add %color1, %color2
+ %7:vec4<f32> = mul %position, %6
+ exit_if # if_1
+ }
+ }
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Inputs_UsedByHelper) {
+ auto* front_facing = b.Var("front_facing", ty.ptr(core::AddressSpace::kIn, ty.bool_()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kFrontFacing;
+ front_facing->SetAttributes(std::move(attributes));
+ }
+ auto* position = b.Var("position", ty.ptr(core::AddressSpace::kIn, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kPosition;
+ attributes.invariant = true;
+ position->SetAttributes(std::move(attributes));
+ }
+ auto* color1 = b.Var("color1", ty.ptr(core::AddressSpace::kIn, ty.f32()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 0;
+ color1->SetAttributes(std::move(attributes));
+ }
+ auto* color2 = b.Var("color2", ty.ptr(core::AddressSpace::kIn, ty.f32()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1;
+ attributes.interpolation = core::Interpolation{core::InterpolationType::kLinear,
+ core::InterpolationSampling::kSample};
+ color2->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(front_facing);
+ mod.root_block->Append(position);
+ mod.root_block->Append(color1);
+ mod.root_block->Append(color2);
+
+ // Inner function has an existing parameter.
+ auto* param = b.FunctionParam("existing_param", ty.f32());
+ auto* foo = b.Function("foo", ty.void_());
+ foo->SetParams({param});
+ b.Append(foo->Block(), [&] {
+ auto* ifelse = b.If(b.Load(front_facing));
+ b.Append(ifelse->True(), [&] {
+ auto* position_value = b.Load(position);
+ auto* color1_value = b.Load(color1);
+ auto* color2_value = b.Load(color2);
+ auto* add = b.Add(ty.f32(), color1_value, color2_value);
+ auto* mul = b.Multiply(ty.vec4<f32>(), position_value, add);
+ b.Divide(ty.vec4<f32>(), mul, param);
+ b.ExitIf(ifelse);
+ });
+ b.Return(foo);
+ });
+
+ // Intermediate function has no existing parameters.
+ auto* bar = b.Function("bar", ty.void_());
+ b.Append(bar->Block(), [&] {
+ b.Call(foo, 42_f);
+ b.Return(bar);
+ });
+
+ auto* ep = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] {
+ b.Call(bar);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %front_facing:ptr<__in, bool, read> = var @builtin(front_facing)
+ %position:ptr<__in, vec4<f32>, read> = var @invariant @builtin(position)
+ %color1:ptr<__in, f32, read> = var @location(0)
+ %color2:ptr<__in, f32, read> = var @location(1) @interpolate(linear, sample)
+}
+
+%foo = func(%existing_param:f32):void {
+ $B2: {
+ %7:bool = load %front_facing
+ if %7 [t: $B3] { # if_1
+ $B3: { # true
+ %8:vec4<f32> = load %position
+ %9:f32 = load %color1
+ %10:f32 = load %color2
+ %11:f32 = add %9, %10
+ %12:vec4<f32> = mul %8, %11
+ %13:vec4<f32> = div %12, %existing_param
+ exit_if # if_1
+ }
+ }
+ ret
+ }
+}
+%bar = func():void {
+ $B4: {
+ %15:void = call %foo, 42.0f
+ ret
+ }
+}
+%main = @fragment func():void {
+ $B5: {
+ %17:void = call %bar
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%existing_param:f32, %front_facing:bool, %position:vec4<f32>, %color1:f32, %color2:f32):void {
+ $B1: {
+ if %front_facing [t: $B2] { # if_1
+ $B2: { # true
+ %7:f32 = add %color1, %color2
+ %8:vec4<f32> = mul %position, %7
+ %9:vec4<f32> = div %8, %existing_param
+ exit_if # if_1
+ }
+ }
+ ret
+ }
+}
+%bar = func(%front_facing_1:bool, %position_1:vec4<f32>, %color1_1:f32, %color2_1:f32):void { # %front_facing_1: 'front_facing', %position_1: 'position', %color1_1: 'color1', %color2_1: 'color2'
+ $B3: {
+ %15:void = call %foo, 42.0f, %front_facing_1, %position_1, %color1_1, %color2_1
+ ret
+ }
+}
+%main = @fragment func(%front_facing_2:bool [@front_facing], %position_2:vec4<f32> [@invariant, @position], %color1_2:f32 [@location(0)], %color2_2:f32 [@location(1), @interpolate(linear, sample)]):void { # %front_facing_2: 'front_facing', %position_2: 'position', %color1_2: 'color1', %color2_2: 'color2'
+ $B4: {
+ %21:void = call %bar, %front_facing_2, %position_2, %color1_2, %color2_2
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Inputs_UsedEntryPointAndHelper) {
+ auto* gid = b.Var("gid", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kGlobalInvocationId;
+ gid->SetAttributes(std::move(attributes));
+ }
+ auto* lid = b.Var("lid", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kLocalInvocationId;
+ lid->SetAttributes(std::move(attributes));
+ }
+ auto* group_id = b.Var("group_id", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kWorkgroupId;
+ group_id->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(gid);
+ mod.root_block->Append(lid);
+ mod.root_block->Append(group_id);
+
+ // Use a subset of the inputs in the helper.
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* gid_value = b.Load(gid);
+ auto* lid_value = b.Load(lid);
+ b.Add(ty.vec3<u32>(), gid_value, lid_value);
+ b.Return(foo);
+ });
+
+ // Use a different subset of the inputs in the entry point.
+ auto* ep = b.Function("main1", ty.void_(), core::ir::Function::PipelineStage::kCompute);
+ b.Append(ep->Block(), [&] {
+ auto* group_value = b.Load(group_id);
+ auto* gid_value = b.Load(gid);
+ b.Add(ty.vec3<u32>(), group_value, gid_value);
+ b.Call(foo);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %gid:ptr<__in, vec3<u32>, read> = var @builtin(global_invocation_id)
+ %lid:ptr<__in, vec3<u32>, read> = var @builtin(local_invocation_id)
+ %group_id:ptr<__in, vec3<u32>, read> = var @builtin(workgroup_id)
+}
+
+%foo = func():void {
+ $B2: {
+ %5:vec3<u32> = load %gid
+ %6:vec3<u32> = load %lid
+ %7:vec3<u32> = add %5, %6
+ ret
+ }
+}
+%main1 = @compute func():void {
+ $B3: {
+ %9:vec3<u32> = load %group_id
+ %10:vec3<u32> = load %gid
+ %11:vec3<u32> = add %9, %10
+ %12:void = call %foo
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%gid:vec3<u32>, %lid:vec3<u32>):void {
+ $B1: {
+ %4:vec3<u32> = add %gid, %lid
+ ret
+ }
+}
+%main1 = @compute func(%gid_1:vec3<u32> [@global_invocation_id], %lid_1:vec3<u32> [@local_invocation_id], %group_id:vec3<u32> [@workgroup_id]):void { # %gid_1: 'gid', %lid_1: 'lid'
+ $B2: {
+ %9:vec3<u32> = add %group_id, %gid_1
+ %10:void = call %foo, %gid_1, %lid_1
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Inputs_UsedEntryPointAndHelper_ForwardReference) {
+ auto* gid = b.Var("gid", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kGlobalInvocationId;
+ gid->SetAttributes(std::move(attributes));
+ }
+ auto* lid = b.Var("lid", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kLocalInvocationId;
+ lid->SetAttributes(std::move(attributes));
+ }
+ auto* group_id = b.Var("group_id", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kWorkgroupId;
+ group_id->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(gid);
+ mod.root_block->Append(lid);
+ mod.root_block->Append(group_id);
+
+ auto* ep = b.Function("main1", ty.void_(), core::ir::Function::PipelineStage::kCompute);
+ auto* foo = b.Function("foo", ty.void_());
+
+ // Use a subset of the inputs in the entry point.
+ b.Append(ep->Block(), [&] {
+ auto* group_value = b.Load(group_id);
+ auto* gid_value = b.Load(gid);
+ b.Add(ty.vec3<u32>(), group_value, gid_value);
+ b.Call(foo);
+ b.Return(ep);
+ });
+
+ // Use a different subset of the variables in the helper.
+ b.Append(foo->Block(), [&] {
+ auto* gid_value = b.Load(gid);
+ auto* lid_value = b.Load(lid);
+ b.Add(ty.vec3<u32>(), gid_value, lid_value);
+ b.Return(foo);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %gid:ptr<__in, vec3<u32>, read> = var @builtin(global_invocation_id)
+ %lid:ptr<__in, vec3<u32>, read> = var @builtin(local_invocation_id)
+ %group_id:ptr<__in, vec3<u32>, read> = var @builtin(workgroup_id)
+}
+
+%main1 = @compute func():void {
+ $B2: {
+ %5:vec3<u32> = load %group_id
+ %6:vec3<u32> = load %gid
+ %7:vec3<u32> = add %5, %6
+ %8:void = call %foo
+ ret
+ }
+}
+%foo = func():void {
+ $B3: {
+ %10:vec3<u32> = load %gid
+ %11:vec3<u32> = load %lid
+ %12:vec3<u32> = add %10, %11
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%main1 = @compute func(%gid:vec3<u32> [@global_invocation_id], %lid:vec3<u32> [@local_invocation_id], %group_id:vec3<u32> [@workgroup_id]):void {
+ $B1: {
+ %5:vec3<u32> = add %group_id, %gid
+ %6:void = call %foo, %gid, %lid
+ ret
+ }
+}
+%foo = func(%gid_1:vec3<u32>, %lid_1:vec3<u32>):void { # %gid_1: 'gid', %lid_1: 'lid'
+ $B2: {
+ %10:vec3<u32> = add %gid_1, %lid_1
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Inputs_UsedByMultipleEntryPoints) {
+ auto* gid = b.Var("gid", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kGlobalInvocationId;
+ gid->SetAttributes(std::move(attributes));
+ }
+ auto* lid = b.Var("lid", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kLocalInvocationId;
+ lid->SetAttributes(std::move(attributes));
+ }
+ auto* group_id = b.Var("group_id", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kWorkgroupId;
+ group_id->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(gid);
+ mod.root_block->Append(lid);
+ mod.root_block->Append(group_id);
+
+ // Use a subset of the inputs in the helper.
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* gid_value = b.Load(gid);
+ auto* lid_value = b.Load(lid);
+ b.Add(ty.vec3<u32>(), gid_value, lid_value);
+ b.Return(foo);
+ });
+
+ // Call the helper without directly referencing any inputs.
+ auto* ep1 = b.Function("main1", ty.void_(), core::ir::Function::PipelineStage::kCompute);
+ b.Append(ep1->Block(), [&] {
+ b.Call(foo);
+ b.Return(ep1);
+ });
+
+ // Reference another input and then call the helper.
+ auto* ep2 = b.Function("main2", ty.void_(), core::ir::Function::PipelineStage::kCompute);
+ b.Append(ep2->Block(), [&] {
+ auto* group_value = b.Load(group_id);
+ b.Add(ty.vec3<u32>(), group_value, group_value);
+ b.Call(foo);
+ b.Return(ep1);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %gid:ptr<__in, vec3<u32>, read> = var @builtin(global_invocation_id)
+ %lid:ptr<__in, vec3<u32>, read> = var @builtin(local_invocation_id)
+ %group_id:ptr<__in, vec3<u32>, read> = var @builtin(workgroup_id)
+}
+
+%foo = func():void {
+ $B2: {
+ %5:vec3<u32> = load %gid
+ %6:vec3<u32> = load %lid
+ %7:vec3<u32> = add %5, %6
+ ret
+ }
+}
+%main1 = @compute func():void {
+ $B3: {
+ %9:void = call %foo
+ ret
+ }
+}
+%main2 = @compute func():void {
+ $B4: {
+ %11:vec3<u32> = load %group_id
+ %12:vec3<u32> = add %11, %11
+ %13:void = call %foo
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%gid:vec3<u32>, %lid:vec3<u32>):void {
+ $B1: {
+ %4:vec3<u32> = add %gid, %lid
+ ret
+ }
+}
+%main1 = @compute func(%gid_1:vec3<u32> [@global_invocation_id], %lid_1:vec3<u32> [@local_invocation_id]):void { # %gid_1: 'gid', %lid_1: 'lid'
+ $B2: {
+ %8:void = call %foo, %gid_1, %lid_1
+ ret
+ }
+}
+%main2 = @compute func(%gid_2:vec3<u32> [@global_invocation_id], %lid_2:vec3<u32> [@local_invocation_id], %group_id:vec3<u32> [@workgroup_id]):void { # %gid_2: 'gid', %lid_2: 'lid'
+ $B3: {
+ %13:vec3<u32> = add %group_id, %group_id
+ %14:void = call %foo, %gid_2, %lid_2
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Input_LoadVectorElement) {
+ auto* lid = b.Var("lid", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kLocalInvocationId;
+ lid->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(lid);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute);
+ b.Append(ep->Block(), [&] {
+ b.LoadVectorElement(lid, 2_u);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %lid:ptr<__in, vec3<u32>, read> = var @builtin(local_invocation_id)
+}
+
+%foo = @compute func():void {
+ $B2: {
+ %3:u32 = load_vector_element %lid, 2u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = @compute func(%lid:vec3<u32> [@local_invocation_id]):void {
+ $B1: {
+ %3:u32 = access %lid, 2u
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Input_AccessChains) {
+ auto* lid = b.Var("lid", ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kLocalInvocationId;
+ lid->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(lid);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute);
+ b.Append(ep->Block(), [&] {
+ auto* access_1 = b.Access(ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()), lid);
+ auto* access_2 = b.Access(ty.ptr(core::AddressSpace::kIn, ty.vec3<u32>()), access_1);
+ auto* vec = b.Load(access_2);
+ auto* z = b.LoadVectorElement(access_2, 2_u);
+ b.Multiply<vec3<u32>>(vec, z);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %lid:ptr<__in, vec3<u32>, read> = var @builtin(local_invocation_id)
+}
+
+%foo = @compute func():void {
+ $B2: {
+ %3:ptr<__in, vec3<u32>, read> = access %lid
+ %4:ptr<__in, vec3<u32>, read> = access %3
+ %5:vec3<u32> = load %4
+ %6:u32 = load_vector_element %4, 2u
+ %7:vec3<u32> = mul %5, %6
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = @compute func(%lid:vec3<u32> [@local_invocation_id]):void {
+ $B1: {
+ %3:u32 = access %lid, 2u
+ %4:vec3<u32> = mul %lid, %3
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Inputs_Struct_LocationOnEachMember) {
+ auto* colors_str = ty.Struct(
+ mod.symbols.New("Colors"),
+ Vector{
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color1"),
+ ty.vec4<f32>(),
+ LocationAttrs(1),
+ },
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color2"),
+ ty.vec4<f32>(),
+ LocationAttrs(2u, core::Interpolation{core::InterpolationType::kLinear,
+ core::InterpolationSampling::kCentroid}),
+ },
+ });
+ auto* colors = b.Var("colors", ty.ptr(core::AddressSpace::kIn, colors_str));
+ mod.root_block->Append(colors);
+
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* ptr = ty.ptr(core::AddressSpace::kIn, ty.vec4<f32>());
+ auto* color1_value = b.Load(b.Access(ptr, colors, 0_u));
+ auto* color2_z_value = b.LoadVectorElement(b.Access(ptr, colors, 1_u), 2_u);
+ b.Multiply(ty.vec4<f32>(), color1_value, color2_z_value);
+ b.Return(foo);
+ });
+
+ auto* ep = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] {
+ b.Call(foo);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0), @location(1)
+ color2:vec4<f32> @offset(16), @location(2), @interpolate(linear, centroid)
+}
+
+$B1: { # root
+ %colors:ptr<__in, Colors, read> = var
+}
+
+%foo = func():void {
+ $B2: {
+ %3:ptr<__in, vec4<f32>, read> = access %colors, 0u
+ %4:vec4<f32> = load %3
+ %5:ptr<__in, vec4<f32>, read> = access %colors, 1u
+ %6:f32 = load_vector_element %5, 2u
+ %7:vec4<f32> = mul %4, %6
+ ret
+ }
+}
+%main = @fragment func():void {
+ $B3: {
+ %9:void = call %foo
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0), @location(1)
+ color2:vec4<f32> @offset(16), @location(2), @interpolate(linear, centroid)
+}
+
+%foo = func(%colors:Colors):void {
+ $B1: {
+ %3:vec4<f32> = access %colors, 0u
+ %4:vec4<f32> = access %colors, 1u
+ %5:f32 = access %4, 2u
+ %6:vec4<f32> = mul %3, %5
+ ret
+ }
+}
+%main = @fragment func(%colors_1:Colors):void { # %colors_1: 'colors'
+ $B2: {
+ %9:void = call %foo, %colors_1
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Inputs_Struct_LocationOnVariable) {
+ auto* colors_str =
+ ty.Struct(mod.symbols.New("Colors"),
+ Vector{
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color1"),
+ ty.vec4<f32>(),
+ },
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color2"),
+ ty.vec4<f32>(),
+ core::type::StructMemberAttributes{
+ /* location */ std::nullopt,
+ /* index */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */
+ core::Interpolation{core::InterpolationType::kPerspective,
+ core::InterpolationSampling::kCentroid},
+ /* invariant */ false,
+ },
+ },
+ });
+ auto* colors = b.Var("colors", ty.ptr(core::AddressSpace::kIn, colors_str));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1u;
+ colors->SetAttributes(attributes);
+ }
+ mod.root_block->Append(colors);
+
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* ptr = ty.ptr(core::AddressSpace::kIn, ty.vec4<f32>());
+ auto* color1_value = b.Load(b.Access(ptr, colors, 0_u));
+ auto* color2_z_value = b.LoadVectorElement(b.Access(ptr, colors, 1_u), 2_u);
+ b.Multiply(ty.vec4<f32>(), color1_value, color2_z_value);
+ b.Return(foo);
+ });
+
+ auto* ep = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] {
+ b.Call(foo);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0)
+ color2:vec4<f32> @offset(16), @interpolate(perspective, centroid)
+}
+
+$B1: { # root
+ %colors:ptr<__in, Colors, read> = var @location(1)
+}
+
+%foo = func():void {
+ $B2: {
+ %3:ptr<__in, vec4<f32>, read> = access %colors, 0u
+ %4:vec4<f32> = load %3
+ %5:ptr<__in, vec4<f32>, read> = access %colors, 1u
+ %6:f32 = load_vector_element %5, 2u
+ %7:vec4<f32> = mul %4, %6
+ ret
+ }
+}
+%main = @fragment func():void {
+ $B3: {
+ %9:void = call %foo
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0), @location(1)
+ color2:vec4<f32> @offset(16), @location(2), @interpolate(perspective, centroid)
+}
+
+%foo = func(%colors:Colors):void {
+ $B1: {
+ %3:vec4<f32> = access %colors, 0u
+ %4:vec4<f32> = access %colors, 1u
+ %5:f32 = access %4, 2u
+ %6:vec4<f32> = mul %3, %5
+ ret
+ }
+}
+%main = @fragment func(%colors_1:Colors):void { # %colors_1: 'colors'
+ $B2: {
+ %9:void = call %foo, %colors_1
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Inputs_Struct_InterpolateOnVariable) {
+ auto* colors_str = ty.Struct(
+ mod.symbols.New("Colors"),
+ Vector{
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color1"),
+ ty.vec4<f32>(),
+ LocationAttrs(1),
+ },
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color2"),
+ ty.vec4<f32>(),
+ LocationAttrs(2u, core::Interpolation{core::InterpolationType::kLinear,
+ core::InterpolationSampling::kSample}),
+ },
+ });
+ auto* colors = b.Var("colors", ty.ptr(core::AddressSpace::kIn, colors_str));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.interpolation = core::Interpolation{core::InterpolationType::kPerspective,
+ core::InterpolationSampling::kCentroid};
+ colors->SetAttributes(attributes);
+ }
+ mod.root_block->Append(colors);
+
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* ptr = ty.ptr(core::AddressSpace::kIn, ty.vec4<f32>());
+ auto* color1_value = b.Load(b.Access(ptr, colors, 0_u));
+ auto* color2_z_value = b.LoadVectorElement(b.Access(ptr, colors, 1_u), 2_u);
+ b.Multiply(ty.vec4<f32>(), color1_value, color2_z_value);
+ b.Return(foo);
+ });
+
+ auto* ep = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] {
+ b.Call(foo);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0), @location(1)
+ color2:vec4<f32> @offset(16), @location(2), @interpolate(linear, sample)
+}
+
+$B1: { # root
+ %colors:ptr<__in, Colors, read> = var @interpolate(perspective, centroid)
+}
+
+%foo = func():void {
+ $B2: {
+ %3:ptr<__in, vec4<f32>, read> = access %colors, 0u
+ %4:vec4<f32> = load %3
+ %5:ptr<__in, vec4<f32>, read> = access %colors, 1u
+ %6:f32 = load_vector_element %5, 2u
+ %7:vec4<f32> = mul %4, %6
+ ret
+ }
+}
+%main = @fragment func():void {
+ $B3: {
+ %9:void = call %foo
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0), @location(1), @interpolate(perspective, centroid)
+ color2:vec4<f32> @offset(16), @location(2), @interpolate(linear, sample)
+}
+
+%foo = func(%colors:Colors):void {
+ $B1: {
+ %3:vec4<f32> = access %colors, 0u
+ %4:vec4<f32> = access %colors, 1u
+ %5:f32 = access %4, 2u
+ %6:vec4<f32> = mul %3, %5
+ ret
+ }
+}
+%main = @fragment func(%colors_1:Colors):void { # %colors_1: 'colors'
+ $B2: {
+ %9:void = call %foo, %colors_1
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Inputs_Struct_LoadWholeStruct) {
+ auto* colors_str = ty.Struct(
+ mod.symbols.New("Colors"),
+ Vector{
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color1"),
+ ty.vec4<f32>(),
+ LocationAttrs(1),
+ },
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color2"),
+ ty.vec4<f32>(),
+ LocationAttrs(2u, core::Interpolation{core::InterpolationType::kLinear,
+ core::InterpolationSampling::kCentroid}),
+ },
+ });
+ auto* colors = b.Var("colors", ty.ptr(core::AddressSpace::kIn, colors_str));
+ mod.root_block->Append(colors);
+
+ auto* foo = b.Function("foo", ty.void_());
+ b.Append(foo->Block(), [&] {
+ auto* load = b.Load(colors);
+ auto* color1_value = b.Access<vec4<f32>>(load, 0_u);
+ auto* color2_z_value = b.Access<f32>(load, 1_u, 2_u);
+ b.Multiply(ty.vec4<f32>(), color1_value, color2_z_value);
+ b.Return(foo);
+ });
+
+ auto* ep = b.Function("main", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] {
+ b.Call(foo);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0), @location(1)
+ color2:vec4<f32> @offset(16), @location(2), @interpolate(linear, centroid)
+}
+
+$B1: { # root
+ %colors:ptr<__in, Colors, read> = var
+}
+
+%foo = func():void {
+ $B2: {
+ %3:Colors = load %colors
+ %4:vec4<f32> = access %3, 0u
+ %5:f32 = access %3, 1u, 2u
+ %6:vec4<f32> = mul %4, %5
+ ret
+ }
+}
+%main = @fragment func():void {
+ $B3: {
+ %8:void = call %foo
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0), @location(1)
+ color2:vec4<f32> @offset(16), @location(2), @interpolate(linear, centroid)
+}
+
+%foo = func(%colors:Colors):void {
+ $B1: {
+ %3:vec4<f32> = access %colors, 0u
+ %4:f32 = access %colors, 1u, 2u
+ %5:vec4<f32> = mul %3, %4
+ ret
+ }
+}
+%main = @fragment func(%colors_1:Colors):void { # %colors_1: 'colors'
+ $B2: {
+ %8:void = call %foo, %colors_1
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, SingleOutput_Builtin) {
+ auto* position = b.Var("position", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kPosition;
+ position->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(position);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kVertex);
+ b.Append(ep->Block(), [&] { //
+ b.Store(position, b.Splat<vec4<f32>>(1_f, 4));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %position:ptr<__out, vec4<f32>, read_write> = var @builtin(position)
+}
+
+%foo = @vertex func():void {
+ $B2: {
+ store %position, vec4<f32>(1.0f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+$B1: { # root
+ %position:ptr<private, vec4<f32>, read_write> = var
+}
+
+%foo_inner = func():void {
+ $B2: {
+ store %position, vec4<f32>(1.0f)
+ ret
+ }
+}
+%foo = @vertex func():vec4<f32> [@position] {
+ $B3: {
+ %4:void = call %foo_inner
+ %5:vec4<f32> = load %position
+ ret %5
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, SingleOutput_Builtin_WithInvariant) {
+ auto* position = b.Var("position", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kPosition;
+ attributes.invariant = true;
+ position->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(position);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kVertex);
+ b.Append(ep->Block(), [&] { //
+ b.Store(position, b.Splat<vec4<f32>>(1_f, 4));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %position:ptr<__out, vec4<f32>, read_write> = var @invariant @builtin(position)
+}
+
+%foo = @vertex func():void {
+ $B2: {
+ store %position, vec4<f32>(1.0f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+$B1: { # root
+ %position:ptr<private, vec4<f32>, read_write> = var
+}
+
+%foo_inner = func():void {
+ $B2: {
+ store %position, vec4<f32>(1.0f)
+ ret
+ }
+}
+%foo = @vertex func():vec4<f32> [@invariant, @position] {
+ $B3: {
+ %4:void = call %foo_inner
+ %5:vec4<f32> = load %position
+ ret %5
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, SingleOutput_Location) {
+ auto* color = b.Var("color", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1u;
+ color->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(color);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] { //
+ b.Store(color, b.Splat<vec4<f32>>(1_f, 4));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %color:ptr<__out, vec4<f32>, read_write> = var @location(1)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ store %color, vec4<f32>(1.0f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+$B1: { # root
+ %color:ptr<private, vec4<f32>, read_write> = var
+}
+
+%foo_inner = func():void {
+ $B2: {
+ store %color, vec4<f32>(1.0f)
+ ret
+ }
+}
+%foo = @fragment func():vec4<f32> [@location(1)] {
+ $B3: {
+ %4:void = call %foo_inner
+ %5:vec4<f32> = load %color
+ ret %5
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, SingleOutput_Location_WithInterpolation) {
+ auto* color = b.Var("color", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1u;
+ attributes.interpolation = core::Interpolation{core::InterpolationType::kPerspective,
+ core::InterpolationSampling::kCentroid};
+ color->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(color);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] { //
+ b.Store(color, b.Splat<vec4<f32>>(1_f, 4));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %color:ptr<__out, vec4<f32>, read_write> = var @location(1) @interpolate(perspective, centroid)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ store %color, vec4<f32>(1.0f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+$B1: { # root
+ %color:ptr<private, vec4<f32>, read_write> = var
+}
+
+%foo_inner = func():void {
+ $B2: {
+ store %color, vec4<f32>(1.0f)
+ ret
+ }
+}
+%foo = @fragment func():vec4<f32> [@location(1), @interpolate(perspective, centroid)] {
+ $B3: {
+ %4:void = call %foo_inner
+ %5:vec4<f32> = load %color
+ ret %5
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, MultipleOutputs) {
+ auto* position = b.Var("position", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kPosition;
+ attributes.invariant = true;
+ position->SetAttributes(std::move(attributes));
+ }
+ auto* color1 = b.Var("color1", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1u;
+ color1->SetAttributes(std::move(attributes));
+ }
+ auto* color2 = b.Var("color2", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1u;
+ attributes.interpolation = core::Interpolation{core::InterpolationType::kPerspective,
+ core::InterpolationSampling::kCentroid};
+ color2->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(position);
+ mod.root_block->Append(color1);
+ mod.root_block->Append(color2);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kVertex);
+ b.Append(ep->Block(), [&] { //
+ b.Store(position, b.Splat<vec4<f32>>(1_f, 4));
+ b.Store(color1, b.Splat<vec4<f32>>(0.5_f, 4));
+ b.Store(color2, b.Splat<vec4<f32>>(0.25_f, 4));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %position:ptr<__out, vec4<f32>, read_write> = var @invariant @builtin(position)
+ %color1:ptr<__out, vec4<f32>, read_write> = var @location(1)
+ %color2:ptr<__out, vec4<f32>, read_write> = var @location(1) @interpolate(perspective, centroid)
+}
+
+%foo = @vertex func():void {
+ $B2: {
+ store %position, vec4<f32>(1.0f)
+ store %color1, vec4<f32>(0.5f)
+ store %color2, vec4<f32>(0.25f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+tint_symbol = struct @align(16) {
+ position:vec4<f32> @offset(0), @invariant, @builtin(position)
+ color1:vec4<f32> @offset(16), @location(1)
+ color2:vec4<f32> @offset(32), @location(1), @interpolate(perspective, centroid)
+}
+
+$B1: { # root
+ %position:ptr<private, vec4<f32>, read_write> = var
+ %color1:ptr<private, vec4<f32>, read_write> = var
+ %color2:ptr<private, vec4<f32>, read_write> = var
+}
+
+%foo_inner = func():void {
+ $B2: {
+ store %position, vec4<f32>(1.0f)
+ store %color1, vec4<f32>(0.5f)
+ store %color2, vec4<f32>(0.25f)
+ ret
+ }
+}
+%foo = @vertex func():tint_symbol {
+ $B3: {
+ %6:void = call %foo_inner
+ %7:vec4<f32> = load %position
+ %8:vec4<f32> = load %color1
+ %9:vec4<f32> = load %color2
+ %10:tint_symbol = construct %7, %8, %9
+ ret %10
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Outputs_Struct_LocationOnEachMember) {
+ auto* builtin_str =
+ ty.Struct(mod.symbols.New("Builtins"), Vector{
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("position"),
+ ty.vec4<f32>(),
+ BuiltinAttrs(core::BuiltinValue::kPosition),
+ },
+ });
+ auto* colors_str = ty.Struct(
+ mod.symbols.New("Colors"),
+ Vector{
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color1"),
+ ty.vec4<f32>(),
+ LocationAttrs(1),
+ },
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color2"),
+ ty.vec4<f32>(),
+ LocationAttrs(2u, core::Interpolation{core::InterpolationType::kPerspective,
+ core::InterpolationSampling::kCentroid}),
+ },
+ });
+
+ auto* builtins = b.Var("builtins", ty.ptr(core::AddressSpace::kOut, builtin_str));
+ auto* colors = b.Var("colors", ty.ptr(core::AddressSpace::kOut, colors_str));
+ mod.root_block->Append(builtins);
+ mod.root_block->Append(colors);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kVertex);
+ b.Append(ep->Block(), [&] { //
+ auto* ptr = ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>());
+ b.Store(b.Access(ptr, builtins, 0_u), b.Splat<vec4<f32>>(1_f, 4));
+ b.Store(b.Access(ptr, colors, 0_u), b.Splat<vec4<f32>>(0.5_f, 4));
+ b.Store(b.Access(ptr, colors, 1_u), b.Splat<vec4<f32>>(0.25_f, 4));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+Builtins = struct @align(16) {
+ position:vec4<f32> @offset(0), @builtin(position)
+}
+
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0), @location(1)
+ color2:vec4<f32> @offset(16), @location(2), @interpolate(perspective, centroid)
+}
+
+$B1: { # root
+ %builtins:ptr<__out, Builtins, read_write> = var
+ %colors:ptr<__out, Colors, read_write> = var
+}
+
+%foo = @vertex func():void {
+ $B2: {
+ %4:ptr<__out, vec4<f32>, read_write> = access %builtins, 0u
+ store %4, vec4<f32>(1.0f)
+ %5:ptr<__out, vec4<f32>, read_write> = access %colors, 0u
+ store %5, vec4<f32>(0.5f)
+ %6:ptr<__out, vec4<f32>, read_write> = access %colors, 1u
+ store %6, vec4<f32>(0.25f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Builtins = struct @align(16) {
+ position:vec4<f32> @offset(0)
+}
+
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0)
+ color2:vec4<f32> @offset(16)
+}
+
+tint_symbol = struct @align(16) {
+ position:vec4<f32> @offset(0), @builtin(position)
+ color1:vec4<f32> @offset(16), @location(1)
+ color2:vec4<f32> @offset(32), @location(2), @interpolate(perspective, centroid)
+}
+
+$B1: { # root
+ %builtins:ptr<private, Builtins, read_write> = var
+ %colors:ptr<private, Colors, read_write> = var
+}
+
+%foo_inner = func():void {
+ $B2: {
+ %4:ptr<private, vec4<f32>, read_write> = access %builtins, 0u
+ store %4, vec4<f32>(1.0f)
+ %5:ptr<private, vec4<f32>, read_write> = access %colors, 0u
+ store %5, vec4<f32>(0.5f)
+ %6:ptr<private, vec4<f32>, read_write> = access %colors, 1u
+ store %6, vec4<f32>(0.25f)
+ ret
+ }
+}
+%foo = @vertex func():tint_symbol {
+ $B3: {
+ %8:void = call %foo_inner
+ %9:ptr<private, vec4<f32>, read_write> = access %builtins, 0u
+ %10:vec4<f32> = load %9
+ %11:ptr<private, vec4<f32>, read_write> = access %colors, 0u
+ %12:vec4<f32> = load %11
+ %13:ptr<private, vec4<f32>, read_write> = access %colors, 1u
+ %14:vec4<f32> = load %13
+ %15:tint_symbol = construct %10, %12, %14
+ ret %15
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Outputs_Struct_LocationOnVariable) {
+ auto* builtin_str =
+ ty.Struct(mod.symbols.New("Builtins"), Vector{
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("position"),
+ ty.vec4<f32>(),
+ BuiltinAttrs(core::BuiltinValue::kPosition),
+ },
+ });
+ auto* colors_str =
+ ty.Struct(mod.symbols.New("Colors"),
+ Vector{
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color1"),
+ ty.vec4<f32>(),
+ },
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color2"),
+ ty.vec4<f32>(),
+ core::type::StructMemberAttributes{
+ /* location */ std::nullopt,
+ /* index */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */
+ core::Interpolation{core::InterpolationType::kPerspective,
+ core::InterpolationSampling::kCentroid},
+ /* invariant */ false,
+ },
+ },
+ });
+
+ auto* builtins = b.Var("builtins", ty.ptr(core::AddressSpace::kOut, builtin_str));
+ auto* colors = b.Var("colors", ty.ptr(core::AddressSpace::kOut, colors_str));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1u;
+ colors->SetAttributes(attributes);
+ }
+ mod.root_block->Append(builtins);
+ mod.root_block->Append(colors);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kVertex);
+ b.Append(ep->Block(), [&] { //
+ auto* ptr = ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>());
+ b.Store(b.Access(ptr, builtins, 0_u), b.Splat<vec4<f32>>(1_f, 4));
+ b.Store(b.Access(ptr, colors, 0_u), b.Splat<vec4<f32>>(0.5_f, 4));
+ b.Store(b.Access(ptr, colors, 1_u), b.Splat<vec4<f32>>(0.25_f, 4));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+Builtins = struct @align(16) {
+ position:vec4<f32> @offset(0), @builtin(position)
+}
+
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0)
+ color2:vec4<f32> @offset(16), @interpolate(perspective, centroid)
+}
+
+$B1: { # root
+ %builtins:ptr<__out, Builtins, read_write> = var
+ %colors:ptr<__out, Colors, read_write> = var @location(1)
+}
+
+%foo = @vertex func():void {
+ $B2: {
+ %4:ptr<__out, vec4<f32>, read_write> = access %builtins, 0u
+ store %4, vec4<f32>(1.0f)
+ %5:ptr<__out, vec4<f32>, read_write> = access %colors, 0u
+ store %5, vec4<f32>(0.5f)
+ %6:ptr<__out, vec4<f32>, read_write> = access %colors, 1u
+ store %6, vec4<f32>(0.25f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Builtins = struct @align(16) {
+ position:vec4<f32> @offset(0)
+}
+
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0)
+ color2:vec4<f32> @offset(16)
+}
+
+tint_symbol = struct @align(16) {
+ position:vec4<f32> @offset(0), @builtin(position)
+ color1:vec4<f32> @offset(16), @location(1)
+ color2:vec4<f32> @offset(32), @location(2), @interpolate(perspective, centroid)
+}
+
+$B1: { # root
+ %builtins:ptr<private, Builtins, read_write> = var
+ %colors:ptr<private, Colors, read_write> = var
+}
+
+%foo_inner = func():void {
+ $B2: {
+ %4:ptr<private, vec4<f32>, read_write> = access %builtins, 0u
+ store %4, vec4<f32>(1.0f)
+ %5:ptr<private, vec4<f32>, read_write> = access %colors, 0u
+ store %5, vec4<f32>(0.5f)
+ %6:ptr<private, vec4<f32>, read_write> = access %colors, 1u
+ store %6, vec4<f32>(0.25f)
+ ret
+ }
+}
+%foo = @vertex func():tint_symbol {
+ $B3: {
+ %8:void = call %foo_inner
+ %9:ptr<private, vec4<f32>, read_write> = access %builtins, 0u
+ %10:vec4<f32> = load %9
+ %11:ptr<private, vec4<f32>, read_write> = access %colors, 0u
+ %12:vec4<f32> = load %11
+ %13:ptr<private, vec4<f32>, read_write> = access %colors, 1u
+ %14:vec4<f32> = load %13
+ %15:tint_symbol = construct %10, %12, %14
+ ret %15
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Outputs_Struct_InterpolateOnVariable) {
+ auto* builtin_str =
+ ty.Struct(mod.symbols.New("Builtins"), Vector{
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("position"),
+ ty.vec4<f32>(),
+ BuiltinAttrs(core::BuiltinValue::kPosition),
+ },
+ });
+ auto* colors_str =
+ ty.Struct(mod.symbols.New("Colors"),
+ Vector{
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color1"),
+ ty.vec4<f32>(),
+ LocationAttrs(2),
+ },
+ core::type::Manager::StructMemberDesc{
+ mod.symbols.New("color2"),
+ ty.vec4<f32>(),
+ LocationAttrs(3, core::Interpolation{core::InterpolationType::kFlat}),
+ },
+ });
+
+ auto* builtins = b.Var("builtins", ty.ptr(core::AddressSpace::kOut, builtin_str));
+ auto* colors = b.Var("colors", ty.ptr(core::AddressSpace::kOut, colors_str));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.interpolation = core::Interpolation{core::InterpolationType::kPerspective,
+ core::InterpolationSampling::kCentroid};
+ colors->SetAttributes(attributes);
+ }
+ mod.root_block->Append(builtins);
+ mod.root_block->Append(colors);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kVertex);
+ b.Append(ep->Block(), [&] { //
+ auto* ptr = ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>());
+ b.Store(b.Access(ptr, builtins, 0_u), b.Splat<vec4<f32>>(1_f, 4));
+ b.Store(b.Access(ptr, colors, 0_u), b.Splat<vec4<f32>>(0.5_f, 4));
+ b.Store(b.Access(ptr, colors, 1_u), b.Splat<vec4<f32>>(0.25_f, 4));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+Builtins = struct @align(16) {
+ position:vec4<f32> @offset(0), @builtin(position)
+}
+
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0), @location(2)
+ color2:vec4<f32> @offset(16), @location(3), @interpolate(flat)
+}
+
+$B1: { # root
+ %builtins:ptr<__out, Builtins, read_write> = var
+ %colors:ptr<__out, Colors, read_write> = var @interpolate(perspective, centroid)
+}
+
+%foo = @vertex func():void {
+ $B2: {
+ %4:ptr<__out, vec4<f32>, read_write> = access %builtins, 0u
+ store %4, vec4<f32>(1.0f)
+ %5:ptr<__out, vec4<f32>, read_write> = access %colors, 0u
+ store %5, vec4<f32>(0.5f)
+ %6:ptr<__out, vec4<f32>, read_write> = access %colors, 1u
+ store %6, vec4<f32>(0.25f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Builtins = struct @align(16) {
+ position:vec4<f32> @offset(0)
+}
+
+Colors = struct @align(16) {
+ color1:vec4<f32> @offset(0)
+ color2:vec4<f32> @offset(16)
+}
+
+tint_symbol = struct @align(16) {
+ position:vec4<f32> @offset(0), @builtin(position)
+ color1:vec4<f32> @offset(16), @location(2), @interpolate(perspective, centroid)
+ color2:vec4<f32> @offset(32), @location(3), @interpolate(flat)
+}
+
+$B1: { # root
+ %builtins:ptr<private, Builtins, read_write> = var
+ %colors:ptr<private, Colors, read_write> = var
+}
+
+%foo_inner = func():void {
+ $B2: {
+ %4:ptr<private, vec4<f32>, read_write> = access %builtins, 0u
+ store %4, vec4<f32>(1.0f)
+ %5:ptr<private, vec4<f32>, read_write> = access %colors, 0u
+ store %5, vec4<f32>(0.5f)
+ %6:ptr<private, vec4<f32>, read_write> = access %colors, 1u
+ store %6, vec4<f32>(0.25f)
+ ret
+ }
+}
+%foo = @vertex func():tint_symbol {
+ $B3: {
+ %8:void = call %foo_inner
+ %9:ptr<private, vec4<f32>, read_write> = access %builtins, 0u
+ %10:vec4<f32> = load %9
+ %11:ptr<private, vec4<f32>, read_write> = access %colors, 0u
+ %12:vec4<f32> = load %11
+ %13:ptr<private, vec4<f32>, read_write> = access %colors, 1u
+ %14:vec4<f32> = load %13
+ %15:tint_symbol = construct %10, %12, %14
+ ret %15
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Outputs_UsedByMultipleEntryPoints) {
+ auto* position = b.Var("position", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kPosition;
+ attributes.invariant = true;
+ position->SetAttributes(std::move(attributes));
+ }
+ auto* color1 = b.Var("color1", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1u;
+ color1->SetAttributes(std::move(attributes));
+ }
+ auto* color2 = b.Var("color2", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1u;
+ attributes.interpolation = core::Interpolation{core::InterpolationType::kPerspective,
+ core::InterpolationSampling::kCentroid};
+ color2->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(position);
+ mod.root_block->Append(color1);
+ mod.root_block->Append(color2);
+
+ auto* ep1 = b.Function("main1", ty.void_(), core::ir::Function::PipelineStage::kVertex);
+ b.Append(ep1->Block(), [&] { //
+ b.Store(position, b.Splat<vec4<f32>>(1_f, 4));
+ b.Return(ep1);
+ });
+
+ auto* ep2 = b.Function("main2", ty.void_(), core::ir::Function::PipelineStage::kVertex);
+ b.Append(ep2->Block(), [&] { //
+ b.Store(position, b.Splat<vec4<f32>>(1_f, 4));
+ b.Store(color1, b.Splat<vec4<f32>>(0.5_f, 4));
+ b.Return(ep2);
+ });
+
+ auto* ep3 = b.Function("main3", ty.void_(), core::ir::Function::PipelineStage::kVertex);
+ b.Append(ep3->Block(), [&] { //
+ b.Store(position, b.Splat<vec4<f32>>(1_f, 4));
+ b.Store(color2, b.Splat<vec4<f32>>(0.25_f, 4));
+ b.Return(ep3);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %position:ptr<__out, vec4<f32>, read_write> = var @invariant @builtin(position)
+ %color1:ptr<__out, vec4<f32>, read_write> = var @location(1)
+ %color2:ptr<__out, vec4<f32>, read_write> = var @location(1) @interpolate(perspective, centroid)
+}
+
+%main1 = @vertex func():void {
+ $B2: {
+ store %position, vec4<f32>(1.0f)
+ ret
+ }
+}
+%main2 = @vertex func():void {
+ $B3: {
+ store %position, vec4<f32>(1.0f)
+ store %color1, vec4<f32>(0.5f)
+ ret
+ }
+}
+%main3 = @vertex func():void {
+ $B4: {
+ store %position, vec4<f32>(1.0f)
+ store %color2, vec4<f32>(0.25f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+tint_symbol = struct @align(16) {
+ position:vec4<f32> @offset(0), @invariant, @builtin(position)
+ color1:vec4<f32> @offset(16), @location(1)
+}
+
+tint_symbol_1 = struct @align(16) {
+ position:vec4<f32> @offset(0), @invariant, @builtin(position)
+ color2:vec4<f32> @offset(16), @location(1), @interpolate(perspective, centroid)
+}
+
+$B1: { # root
+ %position:ptr<private, vec4<f32>, read_write> = var
+ %color1:ptr<private, vec4<f32>, read_write> = var
+ %color2:ptr<private, vec4<f32>, read_write> = var
+}
+
+%main1_inner = func():void {
+ $B2: {
+ store %position, vec4<f32>(1.0f)
+ ret
+ }
+}
+%main2_inner = func():void {
+ $B3: {
+ store %position, vec4<f32>(1.0f)
+ store %color1, vec4<f32>(0.5f)
+ ret
+ }
+}
+%main3_inner = func():void {
+ $B4: {
+ store %position, vec4<f32>(1.0f)
+ store %color2, vec4<f32>(0.25f)
+ ret
+ }
+}
+%main1 = @vertex func():vec4<f32> [@invariant, @position] {
+ $B5: {
+ %8:void = call %main1_inner
+ %9:vec4<f32> = load %position
+ ret %9
+ }
+}
+%main2 = @vertex func():tint_symbol {
+ $B6: {
+ %11:void = call %main2_inner
+ %12:vec4<f32> = load %position
+ %13:vec4<f32> = load %color1
+ %14:tint_symbol = construct %12, %13
+ ret %14
+ }
+}
+%main3 = @vertex func():tint_symbol_1 {
+ $B7: {
+ %16:void = call %main3_inner
+ %17:vec4<f32> = load %position
+ %18:vec4<f32> = load %color2
+ %19:tint_symbol_1 = construct %17, %18
+ ret %19
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Output_LoadAndStore) {
+ auto* color = b.Var("color", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1u;
+ color->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(color);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] { //
+ b.Store(color, b.Splat<vec4<f32>>(1_f, 4));
+ auto* load = b.Load(color);
+ auto* mul = b.Multiply<vec4<f32>>(load, 2_f);
+ b.Store(color, mul);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %color:ptr<__out, vec4<f32>, read_write> = var @location(1)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ store %color, vec4<f32>(1.0f)
+ %3:vec4<f32> = load %color
+ %4:vec4<f32> = mul %3, 2.0f
+ store %color, %4
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+$B1: { # root
+ %color:ptr<private, vec4<f32>, read_write> = var
+}
+
+%foo_inner = func():void {
+ $B2: {
+ store %color, vec4<f32>(1.0f)
+ %3:vec4<f32> = load %color
+ %4:vec4<f32> = mul %3, 2.0f
+ store %color, %4
+ ret
+ }
+}
+%foo = @fragment func():vec4<f32> [@location(1)] {
+ $B3: {
+ %6:void = call %foo_inner
+ %7:vec4<f32> = load %color
+ ret %7
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Output_LoadVectorElementAndStoreVectorElement) {
+ auto* color = b.Var("color", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1u;
+ color->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(color);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] { //
+ b.Store(color, b.Splat<vec4<f32>>(1_f, 4));
+ auto* load = b.LoadVectorElement(color, 2_u);
+ auto* mul = b.Multiply<f32>(load, 2_f);
+ b.StoreVectorElement(color, 2_u, mul);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %color:ptr<__out, vec4<f32>, read_write> = var @location(1)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ store %color, vec4<f32>(1.0f)
+ %3:f32 = load_vector_element %color, 2u
+ %4:f32 = mul %3, 2.0f
+ store_vector_element %color, 2u, %4
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+$B1: { # root
+ %color:ptr<private, vec4<f32>, read_write> = var
+}
+
+%foo_inner = func():void {
+ $B2: {
+ store %color, vec4<f32>(1.0f)
+ %3:f32 = load_vector_element %color, 2u
+ %4:f32 = mul %3, 2.0f
+ store_vector_element %color, 2u, %4
+ ret
+ }
+}
+%foo = @fragment func():vec4<f32> [@location(1)] {
+ $B3: {
+ %6:void = call %foo_inner
+ %7:vec4<f32> = load %color
+ ret %7
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Output_AccessChain) {
+ auto* color = b.Var("color", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1u;
+ color->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(color);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] { //
+ auto* access_1 = b.Access(ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()), color);
+ auto* access_2 = b.Access(ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()), access_1);
+ auto* load = b.LoadVectorElement(access_2, 2_u);
+ auto* mul = b.Multiply<vec4<f32>>(b.Splat<vec4<f32>>(1_f, 4), load);
+ b.Store(access_2, mul);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %color:ptr<__out, vec4<f32>, read_write> = var @location(1)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ %3:ptr<__out, vec4<f32>, read_write> = access %color
+ %4:ptr<__out, vec4<f32>, read_write> = access %3
+ %5:f32 = load_vector_element %4, 2u
+ %6:vec4<f32> = mul vec4<f32>(1.0f), %5
+ store %4, %6
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+$B1: { # root
+ %color:ptr<private, vec4<f32>, read_write> = var
+}
+
+%foo_inner = func():void {
+ $B2: {
+ %3:ptr<private, vec4<f32>, read_write> = access %color
+ %4:ptr<private, vec4<f32>, read_write> = access %3
+ %5:f32 = load_vector_element %4, 2u
+ %6:vec4<f32> = mul vec4<f32>(1.0f), %5
+ store %4, %6
+ ret
+ }
+}
+%foo = @fragment func():vec4<f32> [@location(1)] {
+ $B3: {
+ %8:void = call %foo_inner
+ %9:vec4<f32> = load %color
+ ret %9
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_ShaderIOTest, Inputs_And_Outputs) {
+ auto* position = b.Var("position", ty.ptr(core::AddressSpace::kIn, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kPosition;
+ attributes.invariant = true;
+ position->SetAttributes(std::move(attributes));
+ }
+ auto* color_in = b.Var("color_in", ty.ptr(core::AddressSpace::kIn, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 0;
+ color_in->SetAttributes(std::move(attributes));
+ }
+ auto* color_out_1 = b.Var("color_out_1", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 1;
+ color_out_1->SetAttributes(std::move(attributes));
+ }
+ auto* color_out_2 = b.Var("color_out_2", ty.ptr(core::AddressSpace::kOut, ty.vec4<f32>()));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.location = 2;
+ color_out_2->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(position);
+ mod.root_block->Append(color_in);
+ mod.root_block->Append(color_out_1);
+ mod.root_block->Append(color_out_2);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] {
+ auto* position_value = b.Load(position);
+ auto* color_in_value = b.Load(color_in);
+ b.Store(color_out_1, position_value);
+ b.Store(color_out_2, color_in_value);
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %position:ptr<__in, vec4<f32>, read> = var @invariant @builtin(position)
+ %color_in:ptr<__in, vec4<f32>, read> = var @location(0)
+ %color_out_1:ptr<__out, vec4<f32>, read_write> = var @location(1)
+ %color_out_2:ptr<__out, vec4<f32>, read_write> = var @location(2)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ %6:vec4<f32> = load %position
+ %7:vec4<f32> = load %color_in
+ store %color_out_1, %6
+ store %color_out_2, %7
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+tint_symbol = struct @align(16) {
+ color_out_1:vec4<f32> @offset(0), @location(1)
+ color_out_2:vec4<f32> @offset(16), @location(2)
+}
+
+$B1: { # root
+ %color_out_1:ptr<private, vec4<f32>, read_write> = var
+ %color_out_2:ptr<private, vec4<f32>, read_write> = var
+}
+
+%foo_inner = func(%position:vec4<f32>, %color_in:vec4<f32>):void {
+ $B2: {
+ store %color_out_1, %position
+ store %color_out_2, %color_in
+ ret
+ }
+}
+%foo = @fragment func(%position_1:vec4<f32> [@invariant, @position], %color_in_1:vec4<f32> [@location(0)]):tint_symbol { # %position_1: 'position', %color_in_1: 'color_in'
+ $B3: {
+ %9:void = call %foo_inner, %position_1, %color_in_1
+ %10:vec4<f32> = load %color_out_1
+ %11:vec4<f32> = load %color_out_2
+ %12:tint_symbol = construct %10, %11
+ ret %12
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+// Test that a sample mask array is converted to a scalar u32 for the entry point.
+TEST_F(SpirvReader_ShaderIOTest, SampleMask) {
+ auto* arr = ty.array<u32, 1>();
+ auto* mask_in = b.Var("mask_in", ty.ptr(core::AddressSpace::kIn, arr));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kSampleMask;
+ mask_in->SetAttributes(std::move(attributes));
+ }
+ auto* mask_out = b.Var("mask_out", ty.ptr(core::AddressSpace::kOut, arr));
+ {
+ core::ir::IOAttributes attributes;
+ attributes.builtin = core::BuiltinValue::kSampleMask;
+ mask_out->SetAttributes(std::move(attributes));
+ }
+ mod.root_block->Append(mask_in);
+ mod.root_block->Append(mask_out);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] {
+ auto* mask_value = b.Load(mask_in);
+ auto* doubled = b.Multiply(ty.u32(), b.Access(ty.u32(), mask_value, 0_u), 2_u);
+ b.Store(mask_out, b.Construct(arr, doubled));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+$B1: { # root
+ %mask_in:ptr<__in, array<u32, 1>, read> = var @builtin(sample_mask)
+ %mask_out:ptr<__out, array<u32, 1>, read_write> = var @builtin(sample_mask)
+}
+
+%foo = @fragment func():void {
+ $B2: {
+ %4:array<u32, 1> = load %mask_in
+ %5:u32 = access %4, 0u
+ %6:u32 = mul %5, 2u
+ %7:array<u32, 1> = construct %6
+ store %mask_out, %7
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+$B1: { # root
+ %mask_out:ptr<private, array<u32, 1>, read_write> = var
+}
+
+%foo_inner = func(%mask_in:array<u32, 1>):void {
+ $B2: {
+ %4:u32 = access %mask_in, 0u
+ %5:u32 = mul %4, 2u
+ %6:array<u32, 1> = construct %5
+ store %mask_out, %6
+ ret
+ }
+}
+%foo = @fragment func(%mask_in_1:u32 [@sample_mask]):u32 [@sample_mask] { # %mask_in_1: 'mask_in'
+ $B3: {
+ %9:array<u32, 1> = construct %mask_in_1
+ %10:void = call %foo_inner, %9
+ %11:array<u32, 1> = load %mask_out
+ %12:u32 = access %11, 0u
+ ret %12
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+} // namespace
+} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/reader_test.cc b/src/tint/lang/spirv/reader/reader_test.cc
index d8ab351..5242819 100644
--- a/src/tint/lang/spirv/reader/reader_test.cc
+++ b/src/tint/lang/spirv/reader/reader_test.cc
@@ -154,5 +154,201 @@
)");
}
+TEST_F(SpirvReaderTest, ShaderInputs) {
+ auto got = Run(R"(
+ OpCapability Shader
+ OpCapability SampleRateShading
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main" %coord %colors
+ OpExecutionMode %main OriginUpperLeft
+ OpDecorate %coord BuiltIn FragCoord
+ OpDecorate %colors Location 1
+ OpMemberDecorate %str 1 NoPerspective
+ %void = OpTypeVoid
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %fn_type = OpTypeFunction %void
+ %str = OpTypeStruct %vec4f %vec4f
+ %u32 = OpTypeInt 32 0
+ %u32_0 = OpConstant %u32 0
+ %u32_1 = OpConstant %u32 1
+
+%_ptr_Input_vec4f = OpTypePointer Input %vec4f
+ %_ptr_Input_str = OpTypePointer Input %str
+ %coord = OpVariable %_ptr_Input_vec4f Input
+ %colors = OpVariable %_ptr_Input_str Input
+
+ %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+ %access_a = OpAccessChain %_ptr_Input_vec4f %colors %u32_0
+ %access_b = OpAccessChain %_ptr_Input_vec4f %colors %u32_1
+ %a = OpLoad %vec4f %access_a
+ %b = OpLoad %vec4f %access_b
+ %c = OpLoad %vec4f %coord
+ %mul = OpFMul %vec4f %a %b
+ %add = OpFAdd %vec4f %mul %c
+ OpReturn
+ OpFunctionEnd
+)");
+ ASSERT_EQ(got, Success);
+ EXPECT_EQ(got, R"(
+tint_symbol_2 = struct @align(16) {
+ tint_symbol:vec4<f32> @offset(0), @location(1)
+ tint_symbol_1:vec4<f32> @offset(16), @location(2), @interpolate(linear, center)
+}
+
+%main = @fragment func(%2:vec4<f32> [@position], %3:tint_symbol_2):void {
+ $B1: {
+ %4:vec4<f32> = access %3, 0u
+ %5:vec4<f32> = access %3, 1u
+ %6:vec4<f32> = mul %4, %5
+ %7:vec4<f32> = add %6, %2
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvReaderTest, ShaderOutputs) {
+ auto got = Run(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main" %depth %colors
+ OpExecutionMode %main OriginUpperLeft
+ OpExecutionMode %main DepthReplacing
+ OpDecorate %depth BuiltIn FragDepth
+ OpDecorate %colors Location 1
+ OpMemberDecorate %str 1 NoPerspective
+ %void = OpTypeVoid
+ %f32 = OpTypeFloat 32
+ %vec4f = OpTypeVector %f32 4
+ %fn_type = OpTypeFunction %void
+ %str = OpTypeStruct %vec4f %vec4f
+ %u32 = OpTypeInt 32 0
+ %u32_0 = OpConstant %u32 0
+ %u32_1 = OpConstant %u32 1
+ %f32_42 = OpConstant %f32 42.0
+ %f32_n1 = OpConstant %f32 -1.0
+ %f32_v4_a = OpConstantComposite %vec4f %f32_42 %f32_42 %f32_42 %f32_n1
+ %f32_v4_b = OpConstantComposite %vec4f %f32_n1 %f32_n1 %f32_n1 %f32_42
+
+%_ptr_Output_f32 = OpTypePointer Output %f32
+%_ptr_Output_vec4f = OpTypePointer Output %vec4f
+ %_ptr_Output_str = OpTypePointer Output %str
+ %depth = OpVariable %_ptr_Output_f32 Output
+ %colors = OpVariable %_ptr_Output_str Output
+
+ %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+ %access_a = OpAccessChain %_ptr_Output_vec4f %colors %u32_0
+ %access_b = OpAccessChain %_ptr_Output_vec4f %colors %u32_1
+ OpStore %access_a %f32_v4_a
+ OpStore %access_b %f32_v4_b
+ OpStore %depth %f32_42
+ OpReturn
+ OpFunctionEnd
+)");
+ ASSERT_EQ(got, Success);
+ EXPECT_EQ(got, R"(
+tint_symbol_2 = struct @align(16) {
+ tint_symbol:vec4<f32> @offset(0)
+ tint_symbol_1:vec4<f32> @offset(16)
+}
+
+tint_symbol_4 = struct @align(16) {
+ tint_symbol_3:f32 @offset(0), @builtin(frag_depth)
+ tint_symbol:vec4<f32> @offset(16), @location(1)
+ tint_symbol_1:vec4<f32> @offset(32), @location(2), @interpolate(linear, center)
+}
+
+$B1: { # root
+ %1:ptr<private, f32, read_write> = var
+ %2:ptr<private, tint_symbol_2, read_write> = var
+}
+
+%main_inner = func():void {
+ $B2: {
+ %4:ptr<private, vec4<f32>, read_write> = access %2, 0u
+ %5:ptr<private, vec4<f32>, read_write> = access %2, 1u
+ store %4, vec4<f32>(42.0f, 42.0f, 42.0f, -1.0f)
+ store %5, vec4<f32>(-1.0f, -1.0f, -1.0f, 42.0f)
+ store %1, 42.0f
+ ret
+ }
+}
+%main = @fragment func():tint_symbol_4 {
+ $B3: {
+ %7:void = call %main_inner
+ %8:f32 = load %1
+ %9:ptr<private, vec4<f32>, read_write> = access %2, 0u
+ %10:vec4<f32> = load %9
+ %11:ptr<private, vec4<f32>, read_write> = access %2, 1u
+ %12:vec4<f32> = load %11
+ %13:tint_symbol_4 = construct %8, %10, %12
+ ret %13
+ }
+}
+)");
+}
+
+TEST_F(SpirvReaderTest, SampleMask) {
+ auto got = Run(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint Fragment %main "main" %mask_in %mask_out
+ OpExecutionMode %main OriginUpperLeft
+ OpDecorate %mask_in BuiltIn SampleMask
+ OpDecorate %mask_out BuiltIn SampleMask
+ %void = OpTypeVoid
+ %fn_type = OpTypeFunction %void
+ %u32 = OpTypeInt 32 0
+ %u32_0 = OpConstant %u32 0
+ %u32_1 = OpConstant %u32 1
+ %arr_u32 = OpTypeArray %u32 %u32_1
+
+%_ptr_Input_u32 = OpTypePointer Input %u32
+%_ptr_Input_arr_u32 = OpTypePointer Input %arr_u32
+%_ptr_Output_u32 = OpTypePointer Output %u32
+%_ptr_Output_arr_u32 = OpTypePointer Output %arr_u32
+ %mask_in = OpVariable %_ptr_Input_arr_u32 Input
+ %mask_out = OpVariable %_ptr_Output_arr_u32 Output
+
+ %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+ %mask_in_0 = OpAccessChain %_ptr_Input_u32 %mask_in %u32_0
+%mask_in_val = OpLoad %u32 %mask_in_0
+ %plus_one = OpIAdd %u32 %mask_in_val %u32_1
+ %mask_out_0 = OpAccessChain %_ptr_Output_u32 %mask_out %u32_0
+ OpStore %mask_out_0 %plus_one
+ OpReturn
+ OpFunctionEnd
+)");
+ ASSERT_EQ(got, Success);
+ EXPECT_EQ(got, R"(
+$B1: { # root
+ %1:ptr<private, array<u32, 1>, read_write> = var
+}
+
+%main_inner = func(%3:array<u32, 1>):void {
+ $B2: {
+ %4:u32 = access %3, 0u
+ %5:u32 = add %4, 1u
+ %6:ptr<private, u32, read_write> = access %1, 0u
+ store %6, %5
+ ret
+ }
+}
+%main = @fragment func(%8:u32 [@sample_mask]):u32 [@sample_mask] {
+ $B3: {
+ %9:array<u32, 1> = construct %8
+ %10:void = call %main_inner, %9
+ %11:array<u32, 1> = load %1
+ %12:u32 = access %11, 0u
+ ret %12
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::spirv::reader
