[hlsl] Add ShaderIO transform.
This CL adds the ShaderIO transform into the HLSL IR backend.
Bug: 42251045
Change-Id: I2bbc00b1d40cae3a0b1227afaf732ee36729499b
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/195475
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
Auto-Submit: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/hlsl/writer/function_test.cc b/src/tint/lang/hlsl/writer/function_test.cc
index e2b8660..ca481a5 100644
--- a/src/tint/lang/hlsl/writer/function_test.cc
+++ b/src/tint/lang/hlsl/writer/function_test.cc
@@ -89,10 +89,14 @@
)");
}
-TEST_F(HlslWriterTest, FunctionEntryPointWithParams) {
+// TODO(dsinclair): Need to pull the struct initializer up to a let.
+TEST_F(HlslWriterTest, DISABLED_FunctionEntryPointWithParams) {
+ core::type::StructMemberAttributes pos_attrs{};
+ pos_attrs.builtin = core::BuiltinValue::kPosition;
+
Vector members{
ty.Get<core::type::StructMember>(b.ir.symbols.New("pos"), ty.vec4<f32>(), 0u, 0u, 16u, 16u,
- core::type::StructMemberAttributes{}),
+ pos_attrs),
};
auto* strct = ty.Struct(b.ir.symbols.New("Interface"), std::move(members));
@@ -107,8 +111,17 @@
float4 pos;
};
+struct main_inputs {
+ float4 Interface_pos : SV_Position;
+};
-void main(Interface p) {
+
+void main_inner(Interface p) {
+}
+
+void main(main_inputs inputs) {
+ Interface v = {float4(inputs.Interface_pos.xyz, 1.0f / inputs.Interface_pos.w)};
+ main_inner(v);
}
)");
@@ -137,7 +150,7 @@
)");
}
-TEST_F(HlslWriterTest, DISABLED_FunctionEntryPointWithInAndOutLocations) {
+TEST_F(HlslWriterTest, FunctionEntryPointWithInAndOutLocations) {
// fn frag_main(@location(0) foo : f32) -> @location(1) f32 {
// return foo;
// }
@@ -146,26 +159,26 @@
foo->SetLocation(0, {});
auto* func = b.Function("frag_main", ty.f32(), core::ir::Function::PipelineStage::kFragment);
+ func->SetParams({foo});
func->SetReturnLocation(1, {});
func->Block()->Append(b.Return(func, foo));
ASSERT_TRUE(Generate()) << err_ << output_.hlsl;
- EXPECT_EQ(output_.hlsl, R"(struct tint_symbol_1 {
+ EXPECT_EQ(output_.hlsl, R"(struct frag_main_outputs {
+ float tint_symbol : SV_Target1;
+};
+
+struct frag_main_inputs {
float foo : TEXCOORD0;
};
-struct tint_symbol_2 {
- float value : SV_Target1;
-};
+
float frag_main_inner(float foo) {
return foo;
}
-tint_symbol_2 frag_main(tint_symbol_1 tint_symbol) {
- float inner_result = frag_main_inner(tint_symbol.foo);
- tint_symbol_2 wrapper_result = (tint_symbol_2)0;
- wrapper_result.value = inner_result;
- return wrapper_result;
+frag_main_outputs frag_main(frag_main_inputs inputs) {
+ return {frag_main_inner(inputs.foo)};
}
)");
diff --git a/src/tint/lang/hlsl/writer/if_test.cc b/src/tint/lang/hlsl/writer/if_test.cc
index 750a44c..a771728 100644
--- a/src/tint/lang/hlsl/writer/if_test.cc
+++ b/src/tint/lang/hlsl/writer/if_test.cc
@@ -192,8 +192,7 @@
}
TEST_F(HlslWriterTest, IfWithMultiPhiReturn1) {
- auto* func = b.Function("foo", ty.i32(), core::ir::Function::PipelineStage::kCompute);
- func->SetWorkgroupSize(1, 1, 1);
+ auto* func = b.Function("foo", ty.i32());
b.Append(func->Block(), [&] {
auto* i = b.If(true);
i->SetResults(b.InstructionResult(ty.i32()), b.InstructionResult(ty.bool_()));
@@ -208,7 +207,6 @@
ASSERT_TRUE(Generate()) << err_ << output_.hlsl;
EXPECT_EQ(output_.hlsl, R"(
-[numthreads(1, 1, 1)]
int foo() {
int v = 0;
bool v_1 = false;
@@ -222,12 +220,15 @@
return v;
}
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+}
+
)");
}
TEST_F(HlslWriterTest, IfWithMultiPhiReturn2) {
- auto* func = b.Function("foo", ty.bool_(), core::ir::Function::PipelineStage::kCompute);
- func->SetWorkgroupSize(1, 1, 1);
+ auto* func = b.Function("foo", ty.bool_());
b.Append(func->Block(), [&] {
auto* i = b.If(true);
i->SetResults(b.InstructionResult(ty.i32()), b.InstructionResult(ty.bool_()));
@@ -242,7 +243,6 @@
ASSERT_TRUE(Generate()) << err_ << output_.hlsl;
EXPECT_EQ(output_.hlsl, R"(
-[numthreads(1, 1, 1)]
bool foo() {
int v = 0;
bool v_1 = false;
@@ -256,6 +256,10 @@
return v_1;
}
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+}
+
)");
}
diff --git a/src/tint/lang/hlsl/writer/raise/BUILD.bazel b/src/tint/lang/hlsl/writer/raise/BUILD.bazel
index b716ce7..32f744f 100644
--- a/src/tint/lang/hlsl/writer/raise/BUILD.bazel
+++ b/src/tint/lang/hlsl/writer/raise/BUILD.bazel
@@ -42,11 +42,13 @@
"builtin_polyfill.cc",
"fxc_polyfill.cc",
"raise.cc",
+ "shader_io.cc",
],
hdrs = [
"builtin_polyfill.h",
"fxc_polyfill.h",
"raise.h",
+ "shader_io.h",
],
deps = [
"//src/tint/api/common",
@@ -83,6 +85,7 @@
srcs = [
"builtin_polyfill_test.cc",
"fxc_polyfill_test.cc",
+ "shader_io_test.cc",
],
deps = [
"//src/tint/api/common",
diff --git a/src/tint/lang/hlsl/writer/raise/BUILD.cmake b/src/tint/lang/hlsl/writer/raise/BUILD.cmake
index fd8ae66..de695c2 100644
--- a/src/tint/lang/hlsl/writer/raise/BUILD.cmake
+++ b/src/tint/lang/hlsl/writer/raise/BUILD.cmake
@@ -45,6 +45,8 @@
lang/hlsl/writer/raise/fxc_polyfill.h
lang/hlsl/writer/raise/raise.cc
lang/hlsl/writer/raise/raise.h
+ lang/hlsl/writer/raise/shader_io.cc
+ lang/hlsl/writer/raise/shader_io.h
)
tint_target_add_dependencies(tint_lang_hlsl_writer_raise lib
@@ -81,6 +83,7 @@
tint_add_target(tint_lang_hlsl_writer_raise_test test
lang/hlsl/writer/raise/builtin_polyfill_test.cc
lang/hlsl/writer/raise/fxc_polyfill_test.cc
+ lang/hlsl/writer/raise/shader_io_test.cc
)
tint_target_add_dependencies(tint_lang_hlsl_writer_raise_test test
diff --git a/src/tint/lang/hlsl/writer/raise/BUILD.gn b/src/tint/lang/hlsl/writer/raise/BUILD.gn
index ff211e5..3173de2 100644
--- a/src/tint/lang/hlsl/writer/raise/BUILD.gn
+++ b/src/tint/lang/hlsl/writer/raise/BUILD.gn
@@ -50,6 +50,8 @@
"fxc_polyfill.h",
"raise.cc",
"raise.h",
+ "shader_io.cc",
+ "shader_io.h",
]
deps = [
"${tint_src_dir}/api/common",
@@ -83,6 +85,7 @@
sources = [
"builtin_polyfill_test.cc",
"fxc_polyfill_test.cc",
+ "shader_io_test.cc",
]
deps = [
"${tint_src_dir}:gmock_and_gtest",
diff --git a/src/tint/lang/hlsl/writer/raise/raise.cc b/src/tint/lang/hlsl/writer/raise/raise.cc
index 2c9b47e..c7055e0 100644
--- a/src/tint/lang/hlsl/writer/raise/raise.cc
+++ b/src/tint/lang/hlsl/writer/raise/raise.cc
@@ -33,6 +33,7 @@
#include "src/tint/lang/hlsl/writer/common/options.h"
#include "src/tint/lang/hlsl/writer/raise/builtin_polyfill.h"
#include "src/tint/lang/hlsl/writer/raise/fxc_polyfill.h"
+#include "src/tint/lang/hlsl/writer/raise/shader_io.h"
#include "src/tint/utils/result/result.h"
namespace tint::hlsl::writer {
@@ -59,6 +60,7 @@
RUN_TRANSFORM(raise::FxcPolyfill);
}
+ RUN_TRANSFORM(raise::ShaderIO);
RUN_TRANSFORM(raise::BuiltinPolyfill);
// These transforms need to be run last as various transforms introduce terminator arguments,
diff --git a/src/tint/lang/hlsl/writer/raise/shader_io.cc b/src/tint/lang/hlsl/writer/raise/shader_io.cc
new file mode 100644
index 0000000..c5ecfc5
--- /dev/null
+++ b/src/tint/lang/hlsl/writer/raise/shader_io.cc
@@ -0,0 +1,293 @@
+// 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/hlsl/writer/raise/shader_io.h"
+
+#include <algorithm>
+#include <memory>
+#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/shader_io.h"
+#include "src/tint/lang/core/ir/validator.h"
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+namespace tint::hlsl::writer::raise {
+
+namespace {
+
+/// PIMPL state for the parts of the shader IO transform specific to HLSL.
+/// For HLSL, move all inputs to a struct passed as an entry point parameter, and wrap outputs in
+/// a structure returned by the entry point.
+struct StateImpl : core::ir::transform::ShaderIOBackendState {
+ /// The input parameter
+ core::ir::FunctionParam* input_param = nullptr;
+
+ Vector<uint32_t, 4> input_indices;
+ Vector<uint32_t, 4> output_indices;
+
+ /// The output struct type.
+ core::type::Struct* output_struct = nullptr;
+
+ /// The output values to return from the entry point.
+ Vector<core::ir::Value*, 4> output_values;
+
+ /// Constructor
+ StateImpl(core::ir::Module& mod, core::ir::Function* f) : ShaderIOBackendState(mod, f) {}
+
+ /// Destructor
+ ~StateImpl() override {}
+
+ /// FXC is sensitive to field order in structures, this is used by StructMemberComparator to
+ /// ensure that FXC is happy with the order of emitted fields.
+ uint32_t BuiltinOrder(core::BuiltinValue builtin) {
+ switch (builtin) {
+ case core::BuiltinValue::kPosition:
+ return 1;
+ case core::BuiltinValue::kVertexIndex:
+ return 2;
+ case core::BuiltinValue::kInstanceIndex:
+ return 3;
+ case core::BuiltinValue::kFrontFacing:
+ return 4;
+ case core::BuiltinValue::kFragDepth:
+ return 5;
+ case core::BuiltinValue::kLocalInvocationId:
+ return 6;
+ case core::BuiltinValue::kLocalInvocationIndex:
+ return 7;
+ case core::BuiltinValue::kGlobalInvocationId:
+ return 8;
+ case core::BuiltinValue::kWorkgroupId:
+ return 9;
+ case core::BuiltinValue::kNumWorkgroups:
+ return 10;
+ case core::BuiltinValue::kSampleIndex:
+ return 11;
+ case core::BuiltinValue::kSampleMask:
+ return 12;
+ case core::BuiltinValue::kPointSize:
+ return 13;
+ default:
+ break;
+ }
+ TINT_UNREACHABLE();
+ }
+
+ struct MemberInfo {
+ core::type::Manager::StructMemberDesc member;
+ uint32_t idx;
+ };
+
+ /// Comparison function used to reorder struct members such that all members with
+ /// color attributes appear first (ordered by color slot), then location attributes (ordered by
+ /// location slot), then blend_src attributes (ordered by blend_src slot), followed by those
+ /// with builtin attributes (ordered by BuiltinOrder).
+ /// @param x a struct member
+ /// @param y another struct member
+ /// @returns true if a comes before b
+ bool StructMemberComparator(const MemberInfo& x, const MemberInfo& y) {
+ if (x.member.attributes.color.has_value() && y.member.attributes.color.has_value() &&
+ x.member.attributes.color != y.member.attributes.color) {
+ // Both have color attributes: smallest goes first.
+ return x.member.attributes.color < y.member.attributes.color;
+ } else if (x.member.attributes.color.has_value() != y.member.attributes.color.has_value()) {
+ // The member with the color goes first
+ return x.member.attributes.color.has_value();
+ }
+
+ if (x.member.attributes.location.has_value() && y.member.attributes.location.has_value() &&
+ x.member.attributes.location != y.member.attributes.location) {
+ // Both have location attributes: smallest goes first.
+ return x.member.attributes.location < y.member.attributes.location;
+ } else if (x.member.attributes.location.has_value() !=
+ y.member.attributes.location.has_value()) {
+ // The member with the location goes first
+ return x.member.attributes.location.has_value();
+ }
+
+ if (x.member.attributes.blend_src.has_value() &&
+ y.member.attributes.blend_src.has_value() &&
+ x.member.attributes.blend_src != y.member.attributes.blend_src) {
+ // Both have blend_src attributes: smallest goes first.
+ return x.member.attributes.blend_src < y.member.attributes.blend_src;
+ } else if (x.member.attributes.blend_src.has_value() !=
+ y.member.attributes.blend_src.has_value()) {
+ // The member with the blend_src goes first
+ return x.member.attributes.blend_src.has_value();
+ }
+
+ auto x_blt = x.member.attributes.builtin;
+ auto y_blt = y.member.attributes.builtin;
+ if (x_blt.has_value() && y_blt.has_value()) {
+ // Both are builtins: order matters for FXC.
+ auto order_a = BuiltinOrder(*x_blt);
+ auto order_b = BuiltinOrder(*y_blt);
+ if (order_a != order_b) {
+ return order_a < order_b;
+ }
+ } else if (x_blt.has_value() != y_blt.has_value()) {
+ // The member with the builtin goes first
+ return x_blt.has_value();
+ }
+
+ // Control flow reaches here if x is the same as y.
+ // Sort algorithms sometimes do that.
+ return false;
+ }
+
+ /// @copydoc ShaderIO::BackendState::FinalizeInputs
+ Vector<core::ir::FunctionParam*, 4> FinalizeInputs() override {
+ Vector<core::type::Manager::StructMemberDesc, 4> input_struct_members;
+
+ Vector<MemberInfo, 4> input_data;
+ for (uint32_t i = 0; i < inputs.Length(); ++i) {
+ input_data.Push(MemberInfo{inputs[i], i});
+ }
+
+ input_indices.Resize(inputs.Length());
+
+ // Sort the struct members to satisfy HLSL interfacing matching rules.
+ std::sort(input_data.begin(), input_data.end(),
+ [&](auto& x, auto& y) { return StructMemberComparator(x, y); });
+
+ for (auto input : input_data) {
+ input_indices[input.idx] = static_cast<uint32_t>(input_struct_members.Length());
+ input_struct_members.Push(input.member);
+ }
+
+ if (!input_struct_members.IsEmpty()) {
+ auto* input_struct =
+ ty.Struct(ir.symbols.New(ir.NameOf(func).Name() + "_inputs"), input_struct_members);
+ switch (func->Stage()) {
+ case core::ir::Function::PipelineStage::kFragment:
+ input_struct->AddUsage(core::type::PipelineStageUsage::kFragmentInput);
+ break;
+ case core::ir::Function::PipelineStage::kVertex:
+ input_struct->AddUsage(core::type::PipelineStageUsage::kVertexInput);
+ break;
+ case core::ir::Function::PipelineStage::kCompute:
+ case core::ir::Function::PipelineStage::kUndefined:
+ TINT_UNREACHABLE();
+ }
+ input_param = b.FunctionParam("inputs", input_struct);
+ return {input_param};
+ }
+
+ return tint::Empty;
+ }
+
+ /// @copydoc ShaderIO::BackendState::FinalizeOutputs
+ const core::type::Type* FinalizeOutputs() override {
+ if (outputs.IsEmpty()) {
+ return ty.void_();
+ }
+
+ Vector<MemberInfo, 4> output_data;
+ for (uint32_t i = 0; i < outputs.Length(); ++i) {
+ output_data.Push(MemberInfo{outputs[i], i});
+ }
+
+ // Sort the struct members to satisfy HLSL interfacing matching rules.
+ std::sort(output_data.begin(), output_data.end(),
+ [&](auto& x, auto& y) { return StructMemberComparator(x, y); });
+
+ output_indices.Resize(outputs.Length());
+ output_values.Resize(outputs.Length());
+
+ Vector<core::type::Manager::StructMemberDesc, 4> output_struct_members;
+ for (size_t i = 0; i < output_data.Length(); ++i) {
+ output_indices[output_data[i].idx] = static_cast<uint32_t>(i);
+ output_struct_members.Push(output_data[i].member);
+ }
+
+ output_struct =
+ ty.Struct(ir.symbols.New(ir.NameOf(func).Name() + "_outputs"), output_struct_members);
+ switch (func->Stage()) {
+ case core::ir::Function::PipelineStage::kFragment:
+ output_struct->AddUsage(core::type::PipelineStageUsage::kFragmentOutput);
+ break;
+ case core::ir::Function::PipelineStage::kVertex:
+ output_struct->AddUsage(core::type::PipelineStageUsage::kVertexOutput);
+ break;
+ case core::ir::Function::PipelineStage::kCompute:
+ case core::ir::Function::PipelineStage::kUndefined:
+ TINT_UNREACHABLE();
+ }
+ return output_struct;
+ }
+
+ /// @copydoc ShaderIO::BackendState::GetInput
+ core::ir::Value* GetInput(core::ir::Builder& builder, uint32_t idx) override {
+ auto index = input_indices[idx];
+
+ core::ir::Value* v = builder.Access(inputs[idx].type, input_param, u32(index))->Result(0);
+
+ // If this is an input position builtin we need to invert the 'w' component of the vector.
+ if (inputs[idx].attributes.builtin == core::BuiltinValue::kPosition) {
+ auto* w = builder.Access(ty.f32(), v, 3_u);
+ auto* div = builder.Divide(ty.f32(), 1.0_f, w);
+ auto* swizzle = builder.Swizzle(ty.vec3<f32>(), v, {0, 1, 2});
+ v = builder.Construct(ty.vec4<f32>(), swizzle, div)->Results()[0];
+ }
+
+ return v;
+ }
+
+ /// @copydoc ShaderIO::BackendState::SetOutput
+ void SetOutput(core::ir::Builder&, uint32_t idx, core::ir::Value* value) override {
+ auto index = output_indices[idx];
+ output_values[index] = value;
+ }
+
+ /// @copydoc ShaderIO::BackendState::MakeReturnValue
+ core::ir::Value* MakeReturnValue(core::ir::Builder& builder) override {
+ if (!output_struct) {
+ return nullptr;
+ }
+ return builder.Construct(output_struct, std::move(output_values))->Result(0);
+ }
+};
+} // namespace
+
+Result<SuccessType> ShaderIO(core::ir::Module& ir) {
+ auto result = ValidateAndDumpIfNeeded(ir, "ShaderIO transform");
+ if (result != Success) {
+ return result;
+ }
+
+ core::ir::transform::RunShaderIOBase(ir, [&](core::ir::Module& mod, core::ir::Function* func) {
+ return std::make_unique<StateImpl>(mod, func);
+ });
+
+ return Success;
+}
+
+} // namespace tint::hlsl::writer::raise
diff --git a/src/tint/lang/hlsl/writer/raise/shader_io.h b/src/tint/lang/hlsl/writer/raise/shader_io.h
new file mode 100644
index 0000000..8bc3b86
--- /dev/null
+++ b/src/tint/lang/hlsl/writer/raise/shader_io.h
@@ -0,0 +1,47 @@
+// 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_HLSL_WRITER_RAISE_SHADER_IO_H_
+#define SRC_TINT_LANG_HLSL_WRITER_RAISE_SHADER_IO_H_
+
+#include "src/tint/utils/result/result.h"
+
+// Forward declarations.
+namespace tint::core::ir {
+class Module;
+}
+
+namespace tint::hlsl::writer::raise {
+
+/// ShaderIO is a transform that prepares entry point inputs and outputs for HLSL codegen.
+/// @param module the module to transform
+/// @returns success or failure
+Result<SuccessType> ShaderIO(core::ir::Module& module);
+
+} // namespace tint::hlsl::writer::raise
+
+#endif // SRC_TINT_LANG_HLSL_WRITER_RAISE_SHADER_IO_H_
diff --git a/src/tint/lang/hlsl/writer/raise/shader_io_test.cc b/src/tint/lang/hlsl/writer/raise/shader_io_test.cc
new file mode 100644
index 0000000..47f2b88
--- /dev/null
+++ b/src/tint/lang/hlsl/writer/raise/shader_io_test.cc
@@ -0,0 +1,1017 @@
+// 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 <utility>
+
+#include "src/tint/lang/core/ir/transform/helper_test.h"
+#include "src/tint/lang/core/type/struct.h"
+#include "src/tint/lang/hlsl/writer/raise/shader_io.h"
+
+namespace tint::hlsl::writer::raise {
+namespace {
+
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
+
+using HlslWriterTransformTest = core::ir::transform::TransformTest;
+
+TEST_F(HlslWriterTransformTest, ShaderIONoInputsOrOutputs) {
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute);
+ ep->SetWorkgroupSize(1, 1, 1);
+ b.Append(ep->Block(), [&] { b.Return(ep); });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1, 1, 1) func():void {
+ $B1: {
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = src;
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOParameters_NonStruct) {
+ auto* front_facing = b.FunctionParam("front_facing", ty.bool_());
+ front_facing->SetBuiltin(core::BuiltinValue::kFrontFacing);
+
+ auto* position = b.FunctionParam("position", ty.vec4<f32>());
+ position->SetBuiltin(core::BuiltinValue::kPosition);
+ position->SetInvariant(true);
+
+ auto* color1 = b.FunctionParam("color1", ty.f32());
+ color1->SetLocation(0, {});
+
+ auto* color2 = b.FunctionParam("color2", ty.f32());
+ color2->SetLocation(1, core::Interpolation{core::InterpolationType::kLinear,
+ core::InterpolationSampling::kSample});
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ ep->SetParams({front_facing, position, color1, color2});
+
+ b.Append(ep->Block(), [&] {
+ auto* ifelse = b.If(front_facing);
+ b.Append(ifelse->True(), [&] {
+ b.Multiply(ty.vec4<f32>(), position, b.Add(ty.f32(), color1, color2));
+ b.ExitIf(ifelse);
+ });
+ b.Return(ep);
+ });
+
+ auto* src = 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
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+foo_inputs = struct @align(16) {
+ color1:f32 @offset(0), @location(0)
+ color2:f32 @offset(4), @location(1), @interpolate(linear, sample)
+ position:vec4<f32> @offset(16), @invariant, @builtin(position)
+ front_facing:bool @offset(32), @builtin(front_facing)
+}
+
+%foo_inner = func(%front_facing:bool, %position:vec4<f32>, %color1:f32, %color2:f32):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
+ }
+}
+%foo = @fragment func(%inputs:foo_inputs):void {
+ $B3: {
+ %10:bool = access %inputs, 3u
+ %11:vec4<f32> = access %inputs, 2u
+ %12:f32 = access %11, 3u
+ %13:f32 = div 1.0f, %12
+ %14:vec3<f32> = swizzle %11, xyz
+ %15:vec4<f32> = construct %14, %13
+ %16:f32 = access %inputs, 0u
+ %17:f32 = access %inputs, 1u
+ %18:void = call %foo_inner, %10, %15, %16, %17
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOParameters_Struct) {
+ auto* str_ty = ty.Struct(mod.symbols.New("Inputs"),
+ {
+ {
+ mod.symbols.New("front_facing"),
+ ty.bool_(),
+ core::type::StructMemberAttributes{
+ /* location */ std::nullopt,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ core::BuiltinValue::kFrontFacing,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ {
+ mod.symbols.New("position"),
+ ty.vec4<f32>(),
+ core::type::StructMemberAttributes{
+ /* location */ std::nullopt,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ core::BuiltinValue::kPosition,
+ /* interpolation */ std::nullopt,
+ /* invariant */ true,
+ },
+ },
+ {
+ mod.symbols.New("color1"),
+ ty.f32(),
+ core::type::StructMemberAttributes{
+ /* location */ 0u,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ {
+ mod.symbols.New("color2"),
+ ty.f32(),
+ core::type::StructMemberAttributes{
+ /* location */ 1u,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */
+ core::Interpolation{
+ core::InterpolationType::kLinear,
+ core::InterpolationSampling::kSample,
+ },
+ /* invariant */ false,
+ },
+ },
+ });
+
+ auto* str_param = b.FunctionParam("inputs", str_ty);
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ ep->SetParams({str_param});
+
+ b.Append(ep->Block(), [&] {
+ auto* ifelse = b.If(b.Access(ty.bool_(), str_param, 0_i));
+ b.Append(ifelse->True(), [&] {
+ auto* position = b.Access(ty.vec4<f32>(), str_param, 1_i);
+ auto* color1 = b.Access(ty.f32(), str_param, 2_i);
+ auto* color2 = b.Access(ty.f32(), str_param, 3_i);
+ b.Multiply(ty.vec4<f32>(), position, b.Add(ty.f32(), color1, color2));
+ b.ExitIf(ifelse);
+ });
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+Inputs = struct @align(16) {
+ front_facing:bool @offset(0), @builtin(front_facing)
+ position:vec4<f32> @offset(16), @invariant, @builtin(position)
+ color1:f32 @offset(32), @location(0)
+ color2:f32 @offset(36), @location(1), @interpolate(linear, sample)
+}
+
+%foo = @fragment func(%inputs:Inputs):void {
+ $B1: {
+ %3:bool = access %inputs, 0i
+ if %3 [t: $B2] { # if_1
+ $B2: { # true
+ %4:vec4<f32> = access %inputs, 1i
+ %5:f32 = access %inputs, 2i
+ %6:f32 = access %inputs, 3i
+ %7:f32 = add %5, %6
+ %8:vec4<f32> = mul %4, %7
+ exit_if # if_1
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Inputs = struct @align(16) {
+ front_facing:bool @offset(0)
+ position:vec4<f32> @offset(16)
+ color1:f32 @offset(32)
+ color2:f32 @offset(36)
+}
+
+foo_inputs = struct @align(16) {
+ Inputs_color1:f32 @offset(0), @location(0)
+ Inputs_color2:f32 @offset(4), @location(1), @interpolate(linear, sample)
+ Inputs_position:vec4<f32> @offset(16), @invariant, @builtin(position)
+ Inputs_front_facing:bool @offset(32), @builtin(front_facing)
+}
+
+%foo_inner = func(%inputs:Inputs):void {
+ $B1: {
+ %3:bool = access %inputs, 0i
+ if %3 [t: $B2] { # if_1
+ $B2: { # true
+ %4:vec4<f32> = access %inputs, 1i
+ %5:f32 = access %inputs, 2i
+ %6:f32 = access %inputs, 3i
+ %7:f32 = add %5, %6
+ %8:vec4<f32> = mul %4, %7
+ exit_if # if_1
+ }
+ }
+ ret
+ }
+}
+%foo = @fragment func(%inputs_1:foo_inputs):void { # %inputs_1: 'inputs'
+ $B3: {
+ %11:bool = access %inputs_1, 3u
+ %12:vec4<f32> = access %inputs_1, 2u
+ %13:f32 = access %12, 3u
+ %14:f32 = div 1.0f, %13
+ %15:vec3<f32> = swizzle %12, xyz
+ %16:vec4<f32> = construct %15, %14
+ %17:f32 = access %inputs_1, 0u
+ %18:f32 = access %inputs_1, 1u
+ %19:Inputs = construct %11, %16, %17, %18
+ %20:void = call %foo_inner, %19
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOParameters_Mixed) {
+ auto* str_ty = ty.Struct(mod.symbols.New("Inputs"),
+ {
+ {
+ mod.symbols.New("position"),
+ ty.vec4<f32>(),
+ core::type::StructMemberAttributes{
+ /* location */ std::nullopt,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ core::BuiltinValue::kPosition,
+ /* interpolation */ std::nullopt,
+ /* invariant */ true,
+ },
+ },
+ {
+ mod.symbols.New("color1"),
+ ty.f32(),
+ core::type::StructMemberAttributes{
+ /* location */ 0u,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ });
+
+ auto* front_facing = b.FunctionParam("front_facing", ty.bool_());
+ front_facing->SetBuiltin(core::BuiltinValue::kFrontFacing);
+
+ auto* str_param = b.FunctionParam("inputs", str_ty);
+
+ auto* color2 = b.FunctionParam("color2", ty.f32());
+ color2->SetLocation(1, core::Interpolation{core::InterpolationType::kLinear,
+ core::InterpolationSampling::kSample});
+
+ auto* ep = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+ ep->SetParams({front_facing, str_param, color2});
+
+ b.Append(ep->Block(), [&] {
+ auto* ifelse = b.If(front_facing);
+ b.Append(ifelse->True(), [&] {
+ auto* position = b.Access(ty.vec4<f32>(), str_param, 0_i);
+ auto* color1 = b.Access(ty.f32(), str_param, 1_i);
+ b.Multiply(ty.vec4<f32>(), position, b.Add(ty.f32(), color1, color2));
+ b.ExitIf(ifelse);
+ });
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+Inputs = struct @align(16) {
+ position:vec4<f32> @offset(0), @invariant, @builtin(position)
+ color1:f32 @offset(16), @location(0)
+}
+
+%foo = @fragment func(%front_facing:bool [@front_facing], %inputs:Inputs, %color2:f32 [@location(1), @interpolate(linear, sample)]):void {
+ $B1: {
+ if %front_facing [t: $B2] { # if_1
+ $B2: { # true
+ %5:vec4<f32> = access %inputs, 0i
+ %6:f32 = access %inputs, 1i
+ %7:f32 = add %6, %color2
+ %8:vec4<f32> = mul %5, %7
+ exit_if # if_1
+ }
+ }
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Inputs = struct @align(16) {
+ position:vec4<f32> @offset(0)
+ color1:f32 @offset(16)
+}
+
+foo_inputs = struct @align(16) {
+ Inputs_color1:f32 @offset(0), @location(0)
+ color2:f32 @offset(4), @location(1), @interpolate(linear, sample)
+ Inputs_position:vec4<f32> @offset(16), @invariant, @builtin(position)
+ front_facing:bool @offset(32), @builtin(front_facing)
+}
+
+%foo_inner = func(%front_facing:bool, %inputs:Inputs, %color2:f32):void {
+ $B1: {
+ if %front_facing [t: $B2] { # if_1
+ $B2: { # true
+ %5:vec4<f32> = access %inputs, 0i
+ %6:f32 = access %inputs, 1i
+ %7:f32 = add %6, %color2
+ %8:vec4<f32> = mul %5, %7
+ exit_if # if_1
+ }
+ }
+ ret
+ }
+}
+%foo = @fragment func(%inputs_1:foo_inputs):void { # %inputs_1: 'inputs'
+ $B3: {
+ %11:bool = access %inputs_1, 3u
+ %12:vec4<f32> = access %inputs_1, 2u
+ %13:f32 = access %12, 3u
+ %14:f32 = div 1.0f, %13
+ %15:vec3<f32> = swizzle %12, xyz
+ %16:vec4<f32> = construct %15, %14
+ %17:f32 = access %inputs_1, 0u
+ %18:Inputs = construct %16, %17
+ %19:f32 = access %inputs_1, 1u
+ %20:void = call %foo_inner, %11, %18, %19
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOReturnValue_NonStructBuiltin) {
+ auto* ep = b.Function("foo", ty.vec4<f32>(), core::ir::Function::PipelineStage::kVertex);
+ ep->SetReturnBuiltin(core::BuiltinValue::kPosition);
+ ep->SetReturnInvariant(true);
+
+ b.Append(ep->Block(), [&] { b.Return(ep, b.Construct(ty.vec4<f32>(), 0.5_f)); });
+
+ auto* src = R"(
+%foo = @vertex func():vec4<f32> [@invariant, @position] {
+ $B1: {
+ %2:vec4<f32> = construct 0.5f
+ ret %2
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+foo_outputs = struct @align(16) {
+ tint_symbol:vec4<f32> @offset(0), @invariant, @builtin(position)
+}
+
+%foo_inner = func():vec4<f32> {
+ $B1: {
+ %2:vec4<f32> = construct 0.5f
+ ret %2
+ }
+}
+%foo = @vertex func():foo_outputs {
+ $B2: {
+ %4:vec4<f32> = call %foo_inner
+ %5:foo_outputs = construct %4
+ ret %5
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOReturnValue_NonStructLocation) {
+ auto* ep = b.Function("foo", ty.vec4<f32>(), core::ir::Function::PipelineStage::kFragment);
+ ep->SetReturnLocation(1u, {});
+
+ b.Append(ep->Block(), [&] { b.Return(ep, b.Construct(ty.vec4<f32>(), 0.5_f)); });
+
+ auto* src = R"(
+%foo = @fragment func():vec4<f32> [@location(1)] {
+ $B1: {
+ %2:vec4<f32> = construct 0.5f
+ ret %2
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+foo_outputs = struct @align(16) {
+ tint_symbol:vec4<f32> @offset(0), @location(1)
+}
+
+%foo_inner = func():vec4<f32> {
+ $B1: {
+ %2:vec4<f32> = construct 0.5f
+ ret %2
+ }
+}
+%foo = @fragment func():foo_outputs {
+ $B2: {
+ %4:vec4<f32> = call %foo_inner
+ %5:foo_outputs = construct %4
+ ret %5
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOReturnValue_Struct) {
+ auto* str_ty = ty.Struct(mod.symbols.New("Outputs"),
+ {
+ {
+ mod.symbols.New("position"),
+ ty.vec4<f32>(),
+ core::type::StructMemberAttributes{
+ /* location */ std::nullopt,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ core::BuiltinValue::kPosition,
+ /* interpolation */ std::nullopt,
+ /* invariant */ true,
+ },
+ },
+ {
+ mod.symbols.New("color1"),
+ ty.f32(),
+ core::type::StructMemberAttributes{
+ /* location */ 0u,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ {
+ mod.symbols.New("color2"),
+ ty.f32(),
+ core::type::StructMemberAttributes{
+ /* location */ 1u,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */
+ core::Interpolation{
+ core::InterpolationType::kLinear,
+ core::InterpolationSampling::kSample,
+ },
+ /* invariant */ false,
+ },
+ },
+ });
+
+ auto* ep = b.Function("foo", str_ty, core::ir::Function::PipelineStage::kVertex);
+
+ b.Append(ep->Block(), [&] {
+ b.Return(ep, b.Construct(str_ty, b.Construct(ty.vec4<f32>(), 0_f), 0.25_f, 0.75_f));
+ });
+
+ auto* src = R"(
+Outputs = struct @align(16) {
+ position:vec4<f32> @offset(0), @invariant, @builtin(position)
+ color1:f32 @offset(16), @location(0)
+ color2:f32 @offset(20), @location(1), @interpolate(linear, sample)
+}
+
+%foo = @vertex func():Outputs {
+ $B1: {
+ %2:vec4<f32> = construct 0.0f
+ %3:Outputs = construct %2, 0.25f, 0.75f
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Outputs = struct @align(16) {
+ position:vec4<f32> @offset(0)
+ color1:f32 @offset(16)
+ color2:f32 @offset(20)
+}
+
+foo_outputs = struct @align(16) {
+ Outputs_color1:f32 @offset(0), @location(0)
+ Outputs_color2:f32 @offset(4), @location(1), @interpolate(linear, sample)
+ Outputs_position:vec4<f32> @offset(16), @invariant, @builtin(position)
+}
+
+%foo_inner = func():Outputs {
+ $B1: {
+ %2:vec4<f32> = construct 0.0f
+ %3:Outputs = construct %2, 0.25f, 0.75f
+ ret %3
+ }
+}
+%foo = @vertex func():foo_outputs {
+ $B2: {
+ %5:Outputs = call %foo_inner
+ %6:vec4<f32> = access %5, 0u
+ %7:f32 = access %5, 1u
+ %8:f32 = access %5, 2u
+ %9:foo_outputs = construct %7, %8, %6
+ ret %9
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOReturnValue_DualSourceBlending) {
+ auto* str_ty =
+ ty.Struct(mod.symbols.New("Output"), {
+ {
+ mod.symbols.New("color1"),
+ ty.f32(),
+ core::type::StructMemberAttributes{
+ /* location */ 0u,
+ /* blend_src */ 0u,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ {
+ mod.symbols.New("color2"),
+ ty.f32(),
+ core::type::StructMemberAttributes{
+ /* location */ 0u,
+ /* blend_src */ 1u,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ });
+
+ auto* ep = b.Function("foo", str_ty, core::ir::Function::PipelineStage::kFragment);
+ b.Append(ep->Block(), [&] { b.Return(ep, b.Construct(str_ty, 0.25_f, 0.75_f)); });
+
+ auto* src = R"(
+Output = struct @align(4) {
+ color1:f32 @offset(0), @location(0)
+ color2:f32 @offset(4), @location(0)
+}
+
+%foo = @fragment func():Output {
+ $B1: {
+ %2:Output = construct 0.25f, 0.75f
+ ret %2
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Output = struct @align(4) {
+ color1:f32 @offset(0)
+ color2:f32 @offset(4)
+}
+
+foo_outputs = struct @align(4) {
+ Output_color1:f32 @offset(0), @location(0)
+ Output_color2:f32 @offset(4), @location(0)
+}
+
+%foo_inner = func():Output {
+ $B1: {
+ %2:Output = construct 0.25f, 0.75f
+ ret %2
+ }
+}
+%foo = @fragment func():foo_outputs {
+ $B2: {
+ %4:Output = call %foo_inner
+ %5:f32 = access %4, 0u
+ %6:f32 = access %4, 1u
+ %7:foo_outputs = construct %5, %6
+ ret %7
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOStruct_SharedByVertexAndFragment) {
+ auto* str_ty = ty.Struct(mod.symbols.New("Interface"),
+ {
+ {
+ mod.symbols.New("position"),
+ ty.vec4<f32>(),
+ core::type::StructMemberAttributes{
+ /* location */ std::nullopt,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ core::BuiltinValue::kPosition,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ {
+ mod.symbols.New("color"),
+ ty.vec3<f32>(),
+ core::type::StructMemberAttributes{
+ /* location */ 0u,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ });
+
+ // Vertex shader.
+ {
+ auto* ep = b.Function("vert", str_ty, core::ir::Function::PipelineStage::kVertex);
+
+ b.Append(ep->Block(), [&] {
+ auto* position = b.Construct(ty.vec4<f32>(), 0_f);
+ auto* color = b.Construct(ty.vec3<f32>(), 1_f);
+ b.Return(ep, b.Construct(str_ty, position, color));
+ });
+ }
+
+ // Fragment shader.
+ {
+ auto* inputs = b.FunctionParam("inputs", str_ty);
+
+ auto* ep = b.Function("frag", ty.vec4<f32>(), core::ir::Function::PipelineStage::kFragment);
+ ep->SetParams({inputs});
+ ep->SetReturnLocation(0u, {});
+
+ b.Append(ep->Block(), [&] {
+ auto* position = b.Access(ty.vec4<f32>(), inputs, 0_u);
+ auto* color = b.Access(ty.vec3<f32>(), inputs, 1_u);
+ b.Return(ep, b.Add(ty.vec4<f32>(), position, b.Construct(ty.vec4<f32>(), color, 1_f)));
+ });
+ }
+
+ auto* src = R"(
+Interface = struct @align(16) {
+ position:vec4<f32> @offset(0), @builtin(position)
+ color:vec3<f32> @offset(16), @location(0)
+}
+
+%vert = @vertex func():Interface {
+ $B1: {
+ %2:vec4<f32> = construct 0.0f
+ %3:vec3<f32> = construct 1.0f
+ %4:Interface = construct %2, %3
+ ret %4
+ }
+}
+%frag = @fragment func(%inputs:Interface):vec4<f32> [@location(0)] {
+ $B2: {
+ %7:vec4<f32> = access %inputs, 0u
+ %8:vec3<f32> = access %inputs, 1u
+ %9:vec4<f32> = construct %8, 1.0f
+ %10:vec4<f32> = add %7, %9
+ ret %10
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Interface = struct @align(16) {
+ position:vec4<f32> @offset(0)
+ color:vec3<f32> @offset(16)
+}
+
+vert_outputs = struct @align(16) {
+ Interface_color:vec3<f32> @offset(0), @location(0)
+ Interface_position:vec4<f32> @offset(16), @builtin(position)
+}
+
+frag_inputs = struct @align(16) {
+ Interface_color:vec3<f32> @offset(0), @location(0)
+ Interface_position:vec4<f32> @offset(16), @builtin(position)
+}
+
+frag_outputs = struct @align(16) {
+ tint_symbol:vec4<f32> @offset(0), @location(0)
+}
+
+%vert_inner = func():Interface {
+ $B1: {
+ %2:vec4<f32> = construct 0.0f
+ %3:vec3<f32> = construct 1.0f
+ %4:Interface = construct %2, %3
+ ret %4
+ }
+}
+%frag_inner = func(%inputs:Interface):vec4<f32> {
+ $B2: {
+ %7:vec4<f32> = access %inputs, 0u
+ %8:vec3<f32> = access %inputs, 1u
+ %9:vec4<f32> = construct %8, 1.0f
+ %10:vec4<f32> = add %7, %9
+ ret %10
+ }
+}
+%vert = @vertex func():vert_outputs {
+ $B3: {
+ %12:Interface = call %vert_inner
+ %13:vec4<f32> = access %12, 0u
+ %14:vec3<f32> = access %12, 1u
+ %15:vert_outputs = construct %14, %13
+ ret %15
+ }
+}
+%frag = @fragment func(%inputs_1:frag_inputs):frag_outputs { # %inputs_1: 'inputs'
+ $B4: {
+ %18:vec4<f32> = access %inputs_1, 1u
+ %19:f32 = access %18, 3u
+ %20:f32 = div 1.0f, %19
+ %21:vec3<f32> = swizzle %18, xyz
+ %22:vec4<f32> = construct %21, %20
+ %23:vec3<f32> = access %inputs_1, 0u
+ %24:Interface = construct %22, %23
+ %25:vec4<f32> = call %frag_inner, %24
+ %26:frag_outputs = construct %25
+ ret %26
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(HlslWriterTransformTest, ShaderIOStruct_SharedWithBuffer) {
+ auto* str_ty = ty.Struct(mod.symbols.New("Outputs"),
+ {
+ {
+ mod.symbols.New("position"),
+ ty.vec4<f32>(),
+ core::type::StructMemberAttributes{
+ /* location */ std::nullopt,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ core::BuiltinValue::kPosition,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ {
+ mod.symbols.New("color"),
+ ty.vec3<f32>(),
+ core::type::StructMemberAttributes{
+ /* location */ 0u,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ });
+
+ auto* buffer = mod.root_block->Append(b.Var(ty.ptr(storage, str_ty, read)));
+
+ auto* ep = b.Function("vert", str_ty, core::ir::Function::PipelineStage::kVertex);
+
+ b.Append(ep->Block(), [&] { b.Return(ep, b.Load(buffer)); });
+
+ auto* src = R"(
+Outputs = struct @align(16) {
+ position:vec4<f32> @offset(0), @builtin(position)
+ color:vec3<f32> @offset(16), @location(0)
+}
+
+$B1: { # root
+ %1:ptr<storage, Outputs, read> = var
+}
+
+%vert = @vertex func():Outputs {
+ $B2: {
+ %3:Outputs = load %1
+ ret %3
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Outputs = struct @align(16) {
+ position:vec4<f32> @offset(0)
+ color:vec3<f32> @offset(16)
+}
+
+vert_outputs = struct @align(16) {
+ Outputs_color:vec3<f32> @offset(0), @location(0)
+ Outputs_position:vec4<f32> @offset(16), @builtin(position)
+}
+
+$B1: { # root
+ %1:ptr<storage, Outputs, read> = var
+}
+
+%vert_inner = func():Outputs {
+ $B2: {
+ %3:Outputs = load %1
+ ret %3
+ }
+}
+%vert = @vertex func():vert_outputs {
+ $B3: {
+ %5:Outputs = call %vert_inner
+ %6:vec4<f32> = access %5, 0u
+ %7:vec3<f32> = access %5, 1u
+ %8:vert_outputs = construct %7, %6
+ ret %8
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+// Test that IO attributes are stripped from structures that are not used for the shader interface.
+TEST_F(HlslWriterTransformTest, ShaderIOStructWithAttributes_NotUsedForInterface) {
+ auto* vec4f = ty.vec4<f32>();
+ auto* str_ty = ty.Struct(mod.symbols.New("Outputs"),
+ {
+ {
+ mod.symbols.New("position"),
+ vec4f,
+ core::type::StructMemberAttributes{
+ /* location */ std::nullopt,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ core::BuiltinValue::kPosition,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ {
+ mod.symbols.New("color"),
+ vec4f,
+ core::type::StructMemberAttributes{
+ /* location */ 0u,
+ /* blend_src */ std::nullopt,
+ /* color */ std::nullopt,
+ /* builtin */ std::nullopt,
+ /* interpolation */ std::nullopt,
+ /* invariant */ false,
+ },
+ },
+ });
+
+ auto* buffer = mod.root_block->Append(b.Var(ty.ptr(storage, str_ty, read)));
+
+ auto* ep = b.Function("frag", ty.void_(), core::ir::Function::PipelineStage::kFragment);
+
+ b.Append(ep->Block(), [&] {
+ b.Store(buffer, b.Construct(str_ty));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+Outputs = struct @align(16) {
+ position:vec4<f32> @offset(0), @builtin(position)
+ color:vec4<f32> @offset(16), @location(0)
+}
+
+$B1: { # root
+ %1:ptr<storage, Outputs, read> = var
+}
+
+%frag = @fragment func():void {
+ $B2: {
+ %3:Outputs = construct
+ store %1, %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+Outputs = struct @align(16) {
+ position:vec4<f32> @offset(0)
+ color:vec4<f32> @offset(16)
+}
+
+$B1: { # root
+ %1:ptr<storage, Outputs, read> = var
+}
+
+%frag = @fragment func():void {
+ $B2: {
+ %3:Outputs = construct
+ store %1, %3
+ ret
+ }
+}
+)";
+
+ Run(ShaderIO);
+
+ EXPECT_EQ(expect, str());
+}
+
+} // namespace
+} // namespace tint::hlsl::writer::raise
