[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