[msl] Add polyfill for Dot4x8Packed
Use the new 8-bit integer types to polyfill these functions by casting
to vec4<{i,u}8> and performing the calculations manually.
Bug: 42251016
Change-Id: I2b8eb226a2f1aeadfefca3d6adc980d89debb5e3
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/201156
Reviewed-by: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
diff --git a/src/tint/lang/core/ir/transform/remove_terminator_args.cc b/src/tint/lang/core/ir/transform/remove_terminator_args.cc
index 22f801a..330fe61 100644
--- a/src/tint/lang/core/ir/transform/remove_terminator_args.cc
+++ b/src/tint/lang/core/ir/transform/remove_terminator_args.cc
@@ -166,6 +166,7 @@
Result<SuccessType> RemoveTerminatorArgs(Module& ir) {
auto result = ValidateAndDumpIfNeeded(ir, "RemoveTerminatorArgs transform",
core::ir::Capabilities{
+ core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllowPointersInStructures,
core::ir::Capability::kAllowVectorElementPointer,
});
diff --git a/src/tint/lang/core/ir/transform/rename_conflicts.cc b/src/tint/lang/core/ir/transform/rename_conflicts.cc
index 90de8a1..17733e7 100644
--- a/src/tint/lang/core/ir/transform/rename_conflicts.cc
+++ b/src/tint/lang/core/ir/transform/rename_conflicts.cc
@@ -295,6 +295,7 @@
Result<SuccessType> RenameConflicts(core::ir::Module& ir) {
auto result = ValidateAndDumpIfNeeded(ir, "RenameConflicts transform",
core::ir::Capabilities{
+ core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllowPointersInStructures,
core::ir::Capability::kAllowVectorElementPointer,
});
diff --git a/src/tint/lang/core/ir/transform/value_to_let.cc b/src/tint/lang/core/ir/transform/value_to_let.cc
index 120e0d9..630dd32 100644
--- a/src/tint/lang/core/ir/transform/value_to_let.cc
+++ b/src/tint/lang/core/ir/transform/value_to_let.cc
@@ -176,6 +176,7 @@
Result<SuccessType> ValueToLet(Module& ir) {
auto result = ValidateAndDumpIfNeeded(ir, "ValueToLet transform",
core::ir::Capabilities{
+ core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllowPointersInStructures,
core::ir::Capability::kAllowVectorElementPointer,
});
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index c1919f6..2d6d849 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -77,6 +77,7 @@
#include "src/tint/lang/core/type/f16.h"
#include "src/tint/lang/core/type/f32.h"
#include "src/tint/lang/core/type/i32.h"
+#include "src/tint/lang/core/type/i8.h"
#include "src/tint/lang/core/type/matrix.h"
#include "src/tint/lang/core/type/multisampled_texture.h"
#include "src/tint/lang/core/type/pointer.h"
@@ -84,6 +85,7 @@
#include "src/tint/lang/core/type/storage_texture.h"
#include "src/tint/lang/core/type/texture.h"
#include "src/tint/lang/core/type/u32.h"
+#include "src/tint/lang/core/type/u8.h"
#include "src/tint/lang/core/type/vector.h"
#include "src/tint/lang/core/type/void.h"
#include "src/tint/lang/msl/barrier_type.h"
@@ -119,6 +121,7 @@
auto valid =
core::ir::ValidateAndDumpIfNeeded(ir_, "MSL writer",
core::ir::Capabilities{
+ core::ir::Capability::kAllow8BitIntegers,
core::ir::Capability::kAllowPointersInStructures,
});
if (valid != Success) {
@@ -403,7 +406,7 @@
[&](const core::ir::LoadVectorElement*) { /* inlined */ }, //
[&](const core::ir::Swizzle*) { /* inlined */ }, //
[&](const core::ir::Bitcast*) { /* inlined */ }, //
- [&](const core::ir::CoreBinary*) { /* inlined */ }, //
+ [&](const core::ir::Binary*) { /* inlined */ }, //
[&](const core::ir::CoreUnary*) { /* inlined */ }, //
[&](const core::ir::Load*) { /* inlined */ }, //
[&](const core::ir::Construct*) { /* inlined */ }, //
@@ -419,7 +422,7 @@
[&](const core::ir::InstructionResult* r) {
Switch(
r->Instruction(), //
- [&](const core::ir::CoreBinary* b) { EmitBinary(out, b); }, //
+ [&](const core::ir::Binary* b) { EmitBinary(out, b); }, //
[&](const core::ir::CoreUnary* u) { EmitUnary(out, u); }, //
[&](const core::ir::Convert* b) { EmitConvert(out, b); }, //
[&](const core::ir::Let* l) { out << NameOf(l->Result(0)); }, //
@@ -465,7 +468,7 @@
/// Emit a binary instruction
/// @param b the binary instruction
- void EmitBinary(StringStream& out, const core::ir::CoreBinary* b) {
+ void EmitBinary(StringStream& out, const core::ir::Binary* b) {
auto kind = [&] {
switch (b->Op()) {
case core::BinaryOp::kAdd:
@@ -1100,6 +1103,8 @@
[&](const core::type::F16*) { out << "half"; }, //
[&](const core::type::I32*) { out << "int"; }, //
[&](const core::type::U32*) { out << "uint"; }, //
+ [&](const core::type::I8*) { out << "char"; }, //
+ [&](const core::type::U8*) { out << "uchar"; }, //
[&](const core::type::Array* arr) { EmitArrayType(out, arr); },
[&](const core::type::Vector* vec) { EmitVectorType(out, vec); },
[&](const core::type::Matrix* mat) { EmitMatrixType(out, mat); },
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index 41eb46d..4c183f4 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -46,6 +46,7 @@
#include "src/tint/lang/core/type/vector.h"
#include "src/tint/lang/msl/barrier_type.h"
#include "src/tint/lang/msl/builtin_fn.h"
+#include "src/tint/lang/msl/ir/binary.h"
#include "src/tint/lang/msl/ir/builtin_call.h"
#include "src/tint/lang/msl/ir/component.h"
#include "src/tint/lang/msl/ir/member_builtin_call.h"
@@ -78,6 +79,9 @@
/// A map from an integer vector type to a dot polyfill.
Hashmap<const core::type::Vector*, core::ir::Function*, 4> integer_dot_polyfills{};
+ /// A map from an integer type to a packed 8-bit dot polyfill.
+ Hashmap<const core::type::Type*, core::ir::Function*, 2> packed_8bit_integer_dot_polyfills{};
+
/// Process the module.
void Process() {
// Find the builtins that need replacing.
@@ -98,6 +102,8 @@
case core::BuiltinFn::kAtomicXor:
case core::BuiltinFn::kDistance:
case core::BuiltinFn::kDot:
+ case core::BuiltinFn::kDot4I8Packed:
+ case core::BuiltinFn::kDot4U8Packed:
case core::BuiltinFn::kFrexp:
case core::BuiltinFn::kLength:
case core::BuiltinFn::kModf:
@@ -175,6 +181,10 @@
case core::BuiltinFn::kDot:
Dot(builtin);
break;
+ case core::BuiltinFn::kDot4I8Packed:
+ case core::BuiltinFn::kDot4U8Packed:
+ Dot4x8Packed(builtin);
+ break;
case core::BuiltinFn::kFrexp:
Frexp(builtin);
break;
@@ -371,6 +381,51 @@
builtin->Destroy();
}
+ /// Polyfill a packed 8-bit dot product call.
+ /// @param builtin the builtin call instruction
+ void Dot4x8Packed(core::ir::CoreBuiltinCall* builtin) {
+ b.InsertBefore(builtin, [&] {
+ auto* arg0 = builtin->Args()[0];
+ auto* arg1 = builtin->Args()[1];
+ auto* int32 = builtin->Result(0)->Type();
+ auto* int8 = int32->Is<core::type::I32>()
+ ? static_cast<const core::type::Type*>(ty.i8())
+ : static_cast<const core::type::Type*>(ty.u8());
+ // Calls to packed 8-bit dot products are polyfilled by casting to [u]char4, performing
+ // the dot product, and converting the result to a {i,u}32:
+ // uchar4 vec1 = as_type<uchar4>(param_0);
+ // uchar4 vec2 = as_type<uchar4>(param_1);
+ // result = uint(vec1[0] * vec2[0] + vec1[1] * vec2[1]
+ // + vec1[2] * vec2[2] + vec1[3] * vec2[3]);
+ auto* polyfill = packed_8bit_integer_dot_polyfills.GetOrAdd(int32, [&] {
+ auto* lhs_32 = b.FunctionParam("lhs", ty.u32());
+ auto* rhs_32 = b.FunctionParam("rhs", ty.u32());
+ auto* func = b.Function("tint_packed_8bit_dot", int32);
+ func->SetParams({lhs_32, rhs_32});
+ b.Append(func->Block(), [&] {
+ auto* lhs = b.Bitcast(ty.vec4(int8), lhs_32);
+ auto* rhs = b.Bitcast(ty.vec4(int8), rhs_32);
+ core::ir::Value* sum = nullptr;
+ for (uint32_t i = 0; i < 4; i++) {
+ auto* l = b.Access(int8, lhs, u32(i));
+ auto* r = b.Access(int8, rhs, u32(i));
+ auto* mul = b.Binary<ir::Binary>(core::BinaryOp::kMultiply, int8, l, r);
+ if (sum) {
+ auto* add = b.Binary<ir::Binary>(core::BinaryOp::kAdd, int8, sum, mul);
+ sum = add->Result(0);
+ } else {
+ sum = mul->Result(0);
+ }
+ }
+ b.Return(func, b.Convert(int32, sum));
+ });
+ return func;
+ });
+ b.CallWithResult(builtin->DetachResult(), polyfill, arg0, arg1);
+ });
+ builtin->Destroy();
+ }
+
/// Polyfill a frexp call.
/// @param builtin the builtin call instruction
void Frexp(core::ir::CoreBuiltinCall* builtin) {
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
index 5c1a409..d7c7dfe 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -1022,6 +1022,211 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_BuiltinPolyfillTest, Dot4I8Packed) {
+ auto* value0 = b.FunctionParam<u32>("value0");
+ auto* value1 = b.FunctionParam<u32>("value1");
+ auto* func = b.Function("foo", ty.i32());
+ func->SetParams({value0, value1});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<i32>(core::BuiltinFn::kDot4I8Packed, value0, value1);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%value0:u32, %value1:u32):i32 {
+ $B1: {
+ %4:i32 = dot4I8Packed %value0, %value1
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%value0:u32, %value1:u32):i32 {
+ $B1: {
+ %4:i32 = call %tint_packed_8bit_dot, %value0, %value1
+ ret %4
+ }
+}
+%tint_packed_8bit_dot = func(%lhs:u32, %rhs:u32):i32 {
+ $B2: {
+ %8:vec4<i8> = bitcast %lhs
+ %9:vec4<i8> = bitcast %rhs
+ %10:i8 = access %8, 0u
+ %11:i8 = access %9, 0u
+ %12:i8 = mul %10, %11
+ %13:i8 = access %8, 1u
+ %14:i8 = access %9, 1u
+ %15:i8 = mul %13, %14
+ %16:i8 = add %12, %15
+ %17:i8 = access %8, 2u
+ %18:i8 = access %9, 2u
+ %19:i8 = mul %17, %18
+ %20:i8 = add %16, %19
+ %21:i8 = access %8, 3u
+ %22:i8 = access %9, 3u
+ %23:i8 = mul %21, %22
+ %24:i8 = add %20, %23
+ %25:i32 = convert %24
+ ret %25
+ }
+}
+)";
+
+ capabilities.Add(core::ir::Capability::kAllow8BitIntegers);
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, Dot4U8Packed) {
+ auto* value0 = b.FunctionParam<u32>("value0");
+ auto* value1 = b.FunctionParam<u32>("value1");
+ auto* func = b.Function("foo", ty.u32());
+ func->SetParams({value0, value1});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Call<u32>(core::BuiltinFn::kDot4U8Packed, value0, value1);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%value0:u32, %value1:u32):u32 {
+ $B1: {
+ %4:u32 = dot4U8Packed %value0, %value1
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%value0:u32, %value1:u32):u32 {
+ $B1: {
+ %4:u32 = call %tint_packed_8bit_dot, %value0, %value1
+ ret %4
+ }
+}
+%tint_packed_8bit_dot = func(%lhs:u32, %rhs:u32):u32 {
+ $B2: {
+ %8:vec4<u8> = bitcast %lhs
+ %9:vec4<u8> = bitcast %rhs
+ %10:u8 = access %8, 0u
+ %11:u8 = access %9, 0u
+ %12:u8 = mul %10, %11
+ %13:u8 = access %8, 1u
+ %14:u8 = access %9, 1u
+ %15:u8 = mul %13, %14
+ %16:u8 = add %12, %15
+ %17:u8 = access %8, 2u
+ %18:u8 = access %9, 2u
+ %19:u8 = mul %17, %18
+ %20:u8 = add %16, %19
+ %21:u8 = access %8, 3u
+ %22:u8 = access %9, 3u
+ %23:u8 = mul %21, %22
+ %24:u8 = add %20, %23
+ %25:u32 = convert %24
+ ret %25
+ }
+}
+)";
+
+ capabilities.Add(core::ir::Capability::kAllow8BitIntegers);
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, Dot4x8Packed_MultipleCalls) {
+ auto* v = b.FunctionParam<u32>("v");
+ auto* func = b.Function("foo", ty.void_());
+ func->SetParams({v, v});
+ b.Append(func->Block(), [&] {
+ b.Call<i32>(core::BuiltinFn::kDot4I8Packed, v, v);
+ b.Call<i32>(core::BuiltinFn::kDot4I8Packed, v, v);
+ b.Call<u32>(core::BuiltinFn::kDot4U8Packed, v, v);
+ b.Call<u32>(core::BuiltinFn::kDot4U8Packed, v, v);
+ b.Return(func);
+ });
+
+ auto* src = R"(
+%foo = func(%v:u32%v:u32):void {
+ $B1: {
+ %3:i32 = dot4I8Packed %v, %v
+ %4:i32 = dot4I8Packed %v, %v
+ %5:u32 = dot4U8Packed %v, %v
+ %6:u32 = dot4U8Packed %v, %v
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%v:u32%v:u32):void {
+ $B1: {
+ %3:i32 = call %tint_packed_8bit_dot, %v, %v
+ %5:i32 = call %tint_packed_8bit_dot, %v, %v
+ %6:u32 = call %tint_packed_8bit_dot_1, %v, %v
+ %8:u32 = call %tint_packed_8bit_dot_1, %v, %v
+ ret
+ }
+}
+%tint_packed_8bit_dot = func(%lhs:u32, %rhs:u32):i32 {
+ $B2: {
+ %11:vec4<i8> = bitcast %lhs
+ %12:vec4<i8> = bitcast %rhs
+ %13:i8 = access %11, 0u
+ %14:i8 = access %12, 0u
+ %15:i8 = mul %13, %14
+ %16:i8 = access %11, 1u
+ %17:i8 = access %12, 1u
+ %18:i8 = mul %16, %17
+ %19:i8 = add %15, %18
+ %20:i8 = access %11, 2u
+ %21:i8 = access %12, 2u
+ %22:i8 = mul %20, %21
+ %23:i8 = add %19, %22
+ %24:i8 = access %11, 3u
+ %25:i8 = access %12, 3u
+ %26:i8 = mul %24, %25
+ %27:i8 = add %23, %26
+ %28:i32 = convert %27
+ ret %28
+ }
+}
+%tint_packed_8bit_dot_1 = func(%lhs_1:u32, %rhs_1:u32):u32 { # %tint_packed_8bit_dot_1: 'tint_packed_8bit_dot', %lhs_1: 'lhs', %rhs_1: 'rhs'
+ $B3: {
+ %31:vec4<u8> = bitcast %lhs_1
+ %32:vec4<u8> = bitcast %rhs_1
+ %33:u8 = access %31, 0u
+ %34:u8 = access %32, 0u
+ %35:u8 = mul %33, %34
+ %36:u8 = access %31, 1u
+ %37:u8 = access %32, 1u
+ %38:u8 = mul %36, %37
+ %39:u8 = add %35, %38
+ %40:u8 = access %31, 2u
+ %41:u8 = access %32, 2u
+ %42:u8 = mul %40, %41
+ %43:u8 = add %39, %42
+ %44:u8 = access %31, 3u
+ %45:u8 = access %32, 3u
+ %46:u8 = mul %44, %45
+ %47:u8 = add %43, %46
+ %48:u32 = convert %47
+ ret %48
+ }
+}
+)";
+
+ capabilities.Add(core::ir::Capability::kAllow8BitIntegers);
+ Run(BuiltinPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
TEST_F(MslWriter_BuiltinPolyfillTest, Frexp_Scalar) {
auto* value = b.FunctionParam<f32>("value");
auto* func = b.Function("foo", ty.f32());
diff --git a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.ir.msl
index 00ec0fe..c6f910a 100644
--- a/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot4I8Packed/881e62.wgsl.expected.ir.msl
@@ -1,9 +1,51 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: dot4I8Packed
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct tint_module_vars_struct {
+ device int* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ int prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ int VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+int tint_packed_8bit_dot(int lhs, int rhs) {
+ char4 const v = as_type<char4>(lhs);
+ char4 const v_1 = as_type<char4>(rhs);
+ return int(((((v[0u] * v_1[0u]) + (v[1u] * v_1[1u])) + (v[2u] * v_1[2u])) + (v[3u] * v_1[3u])));
+}
+
+int dot4I8Packed_881e62() {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ int res = tint_packed_8bit_dot(arg_0, arg_1);
+ return res;
+}
+
+fragment void fragment_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot4I8Packed_881e62();
+}
+
+kernel void compute_main(device int* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot4I8Packed_881e62();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot4I8Packed_881e62();
+ return out;
+}
+
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v_2 = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v_2.pos, .VertexOutput_prevent_dce=v_2.prevent_dce};
+}
diff --git a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.ir.msl
index e0f4a9e..0b0c145 100644
--- a/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/dot4U8Packed/fbed7b.wgsl.expected.ir.msl
@@ -1,9 +1,51 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
-../../src/tint/lang/msl/writer/printer/printer.cc:989 internal compiler error: TINT_UNREACHABLE unhandled: dot4U8Packed
-********************************************************************
-* The tint shader compiler has encountered an unexpected error. *
-* *
-* Please help us fix this issue by submitting a bug report at *
-* crbug.com/tint with the source program that triggered the bug. *
-********************************************************************
+struct tint_module_vars_struct {
+ device uint* prevent_dce;
+};
+
+struct VertexOutput {
+ float4 pos;
+ uint prevent_dce;
+};
+
+struct vertex_main_outputs {
+ float4 VertexOutput_pos [[position]];
+ uint VertexOutput_prevent_dce [[user(locn0)]] [[flat]];
+};
+
+uint tint_packed_8bit_dot(uint lhs, uint rhs) {
+ uchar4 const v = as_type<uchar4>(lhs);
+ uchar4 const v_1 = as_type<uchar4>(rhs);
+ return uint(((((v[0u] * v_1[0u]) + (v[1u] * v_1[1u])) + (v[2u] * v_1[2u])) + (v[3u] * v_1[3u])));
+}
+
+uint dot4U8Packed_fbed7b() {
+ uint arg_0 = 1u;
+ uint arg_1 = 1u;
+ uint res = tint_packed_8bit_dot(arg_0, arg_1);
+ return res;
+}
+
+fragment void fragment_main(device uint* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot4U8Packed_fbed7b();
+}
+
+kernel void compute_main(device uint* prevent_dce [[buffer(0)]]) {
+ tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.prevent_dce=prevent_dce};
+ (*tint_module_vars.prevent_dce) = dot4U8Packed_fbed7b();
+}
+
+VertexOutput vertex_main_inner() {
+ VertexOutput out = {};
+ out.pos = float4(0.0f);
+ out.prevent_dce = dot4U8Packed_fbed7b();
+ return out;
+}
+
+vertex vertex_main_outputs vertex_main() {
+ VertexOutput const v_2 = vertex_main_inner();
+ return vertex_main_outputs{.VertexOutput_pos=v_2.pos, .VertexOutput_prevent_dce=v_2.prevent_dce};
+}
