[spirv-reader][ir] Correctly handle GLSL 450 SSign
The SPIR-V `SSign` method allows unsigned arguments and return types.
This is not permitted in WGSL. The SPIR-V spec states that the argument
is treated as a signed value, so bitcast the argument/result as needed.
Bug: 42250952
Change-Id: I89bd7ddd3c2f1f81be6b3e4298090e0ed6cd757e
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/221814
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/core/ir/builder.h b/src/tint/lang/core/ir/builder.h
index 59f6050..75a0e2c 100644
--- a/src/tint/lang/core/ir/builder.h
+++ b/src/tint/lang/core/ir/builder.h
@@ -1152,6 +1152,23 @@
/// Creates a builtin call instruction with an existing instruction result
/// @param result the instruction result to use
/// @param func the builtin function to call
+ /// @param explicit_params the explicit params
+ /// @param args the call arguments
+ /// @returns the instruction
+ template <typename KLASS, typename FUNC, typename... ARGS>
+ tint::traits::EnableIf<tint::traits::IsTypeOrDerived<KLASS, ir::BuiltinCall>, KLASS*>
+ CallExplicitWithResult(ir::InstructionResult* result,
+ FUNC func,
+ VectorRef<const core::type::Type*> explicit_params,
+ ARGS&&... args) {
+ auto* inst = ir.CreateInstruction<KLASS>(result, func, Values(std::forward<ARGS>(args)...));
+ inst->SetExplicitTemplateParams(explicit_params);
+ return Append(inst);
+ }
+
+ /// Creates a builtin call instruction with an existing instruction result
+ /// @param result the instruction result to use
+ /// @param func the builtin function to call
/// @param args the call arguments
/// @returns the instruction
template <typename KLASS, typename FUNC, typename... ARGS>
@@ -1164,6 +1181,22 @@
/// Creates a builtin call instruction
/// @param type the return type of the call
/// @param func the builtin function to call
+ /// @param explicit_params the explicit parameters
+ /// @param args the call arguments
+ /// @returns the instruction
+ template <typename KLASS, typename FUNC, typename... ARGS>
+ tint::traits::EnableIf<tint::traits::IsTypeOrDerived<KLASS, ir::BuiltinCall>, KLASS*>
+ CallExplicit(const core::type::Type* type,
+ FUNC func,
+ VectorRef<const core::type::Type*> explicit_params,
+ ARGS&&... args) {
+ return CallExplicitWithResult<KLASS>(InstructionResult(type), func, explicit_params,
+ Values(std::forward<ARGS>(args)...));
+ }
+
+ /// Creates a builtin call instruction
+ /// @param type the return type of the call
+ /// @param func the builtin function to call
/// @param args the call arguments
/// @returns the instruction
template <typename KLASS, typename FUNC, typename... ARGS>
diff --git a/src/tint/lang/core/ir/call.cc b/src/tint/lang/core/ir/call.cc
index a04b408..a91e1f5 100644
--- a/src/tint/lang/core/ir/call.cc
+++ b/src/tint/lang/core/ir/call.cc
@@ -27,8 +27,6 @@
#include "src/tint/lang/core/ir/call.h"
-#include <utility>
-
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Call);
namespace tint::core::ir {
@@ -37,4 +35,20 @@
Call::~Call() = default;
+std::string Call::ExplicitTemplateParamsToStr() const {
+ std::string explicit_params;
+ auto ep = ExplicitTemplateParams();
+ if (!ep.IsEmpty()) {
+ explicit_params += "<";
+ for (size_t i = 0; i < ep.Length(); ++i) {
+ if (i > 0) {
+ explicit_params += ", ";
+ }
+ explicit_params += ep[i]->FriendlyName();
+ }
+ explicit_params += ">";
+ }
+ return explicit_params;
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/call.h b/src/tint/lang/core/ir/call.h
index fa9951f..8c4099a 100644
--- a/src/tint/lang/core/ir/call.h
+++ b/src/tint/lang/core/ir/call.h
@@ -28,7 +28,10 @@
#ifndef SRC_TINT_LANG_CORE_IR_CALL_H_
#define SRC_TINT_LANG_CORE_IR_CALL_H_
+#include <string>
+
#include "src/tint/lang/core/ir/operand_instruction.h"
+#include "src/tint/lang/core/type/type.h"
#include "src/tint/utils/rtti/castable.h"
namespace tint::core::ir {
@@ -41,6 +44,15 @@
/// @returns the offset of the arguments in Operands()
virtual size_t ArgsOperandOffset() const { return 0; }
+ /// Sets the explicit template params for the call
+ void SetExplicitTemplateParams(VectorRef<const core::type::Type*> params) {
+ explicit_template_params_ = params;
+ }
+ /// Retrieves the explicit template params for the call
+ tint::VectorRef<const core::type::Type*> ExplicitTemplateParams() const {
+ return explicit_template_params_;
+ }
+
/// @returns the call arguments
tint::Slice<Value* const> Args() { return operands_.Slice().Offset(ArgsOperandOffset()); }
@@ -60,9 +72,14 @@
/// @returns the side effects for this instruction
Accesses GetSideEffects() const override { return Accesses{Access::kLoad, Access::kStore}; }
+ /// Returns the explicit template parameters as a string.
+ std::string ExplicitTemplateParamsToStr() const;
+
protected:
/// Constructor
explicit Call(Id id);
+
+ Vector<const core::type::Type*, 1> explicit_template_params_;
};
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 5d1a57a..9458b1d 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -2791,7 +2791,8 @@
};
auto builtin = core::intrinsic::LookupFn(context, call->FriendlyName().c_str(), call->FuncId(),
- Empty, args, core::EvaluationStage::kRuntime);
+ call->ExplicitTemplateParams(), args,
+ core::EvaluationStage::kRuntime);
if (builtin != Success) {
AddError(call) << builtin.Failure();
return;
@@ -2827,9 +2828,9 @@
symbols_,
};
- auto result =
- core::intrinsic::LookupMemberFn(context, call->FriendlyName().c_str(), call->FuncId(),
- Empty, std::move(args), core::EvaluationStage::kRuntime);
+ auto result = core::intrinsic::LookupMemberFn(context, call->FriendlyName().c_str(),
+ call->FuncId(), call->ExplicitTemplateParams(),
+ std::move(args), core::EvaluationStage::kRuntime);
if (result != Success) {
AddError(call) << result.Failure();
return;
diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index f2e5cc9..6016b84 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -112,6 +112,8 @@
return "normalize";
case BuiltinFn::kInverse:
return "inverse";
+ case BuiltinFn::kSign:
+ return "sign";
case BuiltinFn::kSdot:
return "sdot";
case BuiltinFn::kUdot:
@@ -167,6 +169,7 @@
case BuiltinFn::kNone:
case BuiltinFn::kNormalize:
case BuiltinFn::kInverse:
+ case BuiltinFn::kSign:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.cc.tmpl b/src/tint/lang/spirv/builtin_fn.cc.tmpl
index 4285e63..7ffcae8 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -74,6 +74,7 @@
case BuiltinFn::kNone:
case BuiltinFn::kNormalize:
case BuiltinFn::kInverse:
+ case BuiltinFn::kSign:
break;
}
return core::ir::Instruction::Accesses{};
diff --git a/src/tint/lang/spirv/builtin_fn.h b/src/tint/lang/spirv/builtin_fn.h
index 24f65d2..55084ef 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -83,6 +83,7 @@
kVectorTimesScalar,
kNormalize,
kInverse,
+ kSign,
kSdot,
kUdot,
kNone,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 6c08d99..30b644d 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -1303,55 +1303,58 @@
/* [104] */ MatcherIndex(21),
/* [105] */ MatcherIndex(0),
/* [106] */ MatcherIndex(4),
- /* [107] */ MatcherIndex(45),
- /* [108] */ MatcherIndex(33),
- /* [109] */ MatcherIndex(9),
- /* [110] */ MatcherIndex(5),
- /* [111] */ MatcherIndex(9),
- /* [112] */ MatcherIndex(2),
- /* [113] */ MatcherIndex(45),
- /* [114] */ MatcherIndex(34),
- /* [115] */ MatcherIndex(10),
- /* [116] */ MatcherIndex(5),
- /* [117] */ MatcherIndex(45),
- /* [118] */ MatcherIndex(35),
- /* [119] */ MatcherIndex(45),
- /* [120] */ MatcherIndex(36),
- /* [121] */ MatcherIndex(9),
- /* [122] */ MatcherIndex(1),
- /* [123] */ MatcherIndex(10),
- /* [124] */ MatcherIndex(1),
- /* [125] */ MatcherIndex(32),
- /* [126] */ MatcherIndex(0),
- /* [127] */ MatcherIndex(9),
- /* [128] */ MatcherIndex(7),
- /* [129] */ MatcherIndex(9),
- /* [130] */ MatcherIndex(3),
- /* [131] */ MatcherIndex(9),
- /* [132] */ MatcherIndex(8),
- /* [133] */ MatcherIndex(10),
- /* [134] */ MatcherIndex(8),
- /* [135] */ MatcherIndex(10),
- /* [136] */ MatcherIndex(2),
- /* [137] */ MatcherIndex(43),
- /* [138] */ MatcherIndex(0),
- /* [139] */ MatcherIndex(9),
- /* [140] */ MatcherIndex(0),
- /* [141] */ MatcherIndex(10),
- /* [142] */ MatcherIndex(0),
- /* [143] */ MatcherIndex(12),
- /* [144] */ MatcherIndex(0),
- /* [145] */ MatcherIndex(16),
- /* [146] */ MatcherIndex(0),
- /* [147] */ MatcherIndex(20),
- /* [148] */ MatcherIndex(0),
- /* [149] */ MatcherIndex(47),
- /* [150] */ MatcherIndex(6),
- /* [151] */ MatcherIndex(46),
- /* [152] */ MatcherIndex(48),
- /* [153] */ MatcherIndex(37),
- /* [154] */ MatcherIndex(50),
- /* [155] */ MatcherIndex(49),
+ /* [107] */ MatcherIndex(21),
+ /* [108] */ MatcherIndex(2),
+ /* [109] */ MatcherIndex(1),
+ /* [110] */ MatcherIndex(45),
+ /* [111] */ MatcherIndex(33),
+ /* [112] */ MatcherIndex(9),
+ /* [113] */ MatcherIndex(5),
+ /* [114] */ MatcherIndex(9),
+ /* [115] */ MatcherIndex(2),
+ /* [116] */ MatcherIndex(45),
+ /* [117] */ MatcherIndex(34),
+ /* [118] */ MatcherIndex(10),
+ /* [119] */ MatcherIndex(5),
+ /* [120] */ MatcherIndex(45),
+ /* [121] */ MatcherIndex(35),
+ /* [122] */ MatcherIndex(45),
+ /* [123] */ MatcherIndex(36),
+ /* [124] */ MatcherIndex(9),
+ /* [125] */ MatcherIndex(1),
+ /* [126] */ MatcherIndex(10),
+ /* [127] */ MatcherIndex(1),
+ /* [128] */ MatcherIndex(32),
+ /* [129] */ MatcherIndex(0),
+ /* [130] */ MatcherIndex(9),
+ /* [131] */ MatcherIndex(7),
+ /* [132] */ MatcherIndex(9),
+ /* [133] */ MatcherIndex(3),
+ /* [134] */ MatcherIndex(9),
+ /* [135] */ MatcherIndex(8),
+ /* [136] */ MatcherIndex(10),
+ /* [137] */ MatcherIndex(8),
+ /* [138] */ MatcherIndex(10),
+ /* [139] */ MatcherIndex(2),
+ /* [140] */ MatcherIndex(43),
+ /* [141] */ MatcherIndex(0),
+ /* [142] */ MatcherIndex(9),
+ /* [143] */ MatcherIndex(0),
+ /* [144] */ MatcherIndex(10),
+ /* [145] */ MatcherIndex(0),
+ /* [146] */ MatcherIndex(12),
+ /* [147] */ MatcherIndex(0),
+ /* [148] */ MatcherIndex(16),
+ /* [149] */ MatcherIndex(0),
+ /* [150] */ MatcherIndex(20),
+ /* [151] */ MatcherIndex(0),
+ /* [152] */ MatcherIndex(47),
+ /* [153] */ MatcherIndex(6),
+ /* [154] */ MatcherIndex(46),
+ /* [155] */ MatcherIndex(48),
+ /* [156] */ MatcherIndex(37),
+ /* [157] */ MatcherIndex(50),
+ /* [158] */ MatcherIndex(49),
};
static_assert(MatcherIndicesIndex::CanIndex(kMatcherIndices),
@@ -1396,7 +1399,7 @@
{
/* [7] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [8] */
@@ -1406,17 +1409,17 @@
{
/* [9] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [10] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [11] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(111),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [12] */
@@ -1426,7 +1429,7 @@
{
/* [13] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [14] */
@@ -1436,17 +1439,17 @@
{
/* [15] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [16] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [17] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(111),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [18] */
@@ -1456,7 +1459,7 @@
{
/* [19] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [20] */
@@ -1466,27 +1469,27 @@
{
/* [21] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [22] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [23] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(135),
+ /* matcher_indices */ MatcherIndicesIndex(138),
},
{
/* [24] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(107),
+ /* matcher_indices */ MatcherIndicesIndex(110),
},
{
/* [25] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [26] */
@@ -1506,17 +1509,17 @@
{
/* [29] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [30] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(113),
+ /* matcher_indices */ MatcherIndicesIndex(116),
},
{
/* [31] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [32] */
@@ -1536,17 +1539,17 @@
{
/* [35] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [36] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(107),
+ /* matcher_indices */ MatcherIndicesIndex(110),
},
{
/* [37] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [38] */
@@ -1561,17 +1564,17 @@
{
/* [40] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(111),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [41] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(113),
+ /* matcher_indices */ MatcherIndicesIndex(116),
},
{
/* [42] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [43] */
@@ -1586,7 +1589,7 @@
{
/* [45] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(111),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [46] */
@@ -1596,7 +1599,7 @@
{
/* [47] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [48] */
@@ -1611,7 +1614,7 @@
{
/* [50] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(127),
+ /* matcher_indices */ MatcherIndicesIndex(130),
},
{
/* [51] */
@@ -1621,7 +1624,7 @@
{
/* [52] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [53] */
@@ -1636,7 +1639,7 @@
{
/* [55] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(129),
+ /* matcher_indices */ MatcherIndicesIndex(132),
},
{
/* [56] */
@@ -1646,7 +1649,7 @@
{
/* [57] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [58] */
@@ -1661,7 +1664,7 @@
{
/* [60] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(111),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [61] */
@@ -1671,7 +1674,7 @@
{
/* [62] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [63] */
@@ -1686,7 +1689,7 @@
{
/* [65] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(111),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [66] */
@@ -1696,7 +1699,7 @@
{
/* [67] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [68] */
@@ -1711,7 +1714,7 @@
{
/* [70] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(135),
+ /* matcher_indices */ MatcherIndicesIndex(138),
},
{
/* [71] */
@@ -1721,7 +1724,7 @@
{
/* [72] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [73] */
@@ -1731,12 +1734,12 @@
{
/* [74] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [75] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(111),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [76] */
@@ -1746,7 +1749,7 @@
{
/* [77] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [78] */
@@ -1756,12 +1759,12 @@
{
/* [79] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [80] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [81] */
@@ -1781,22 +1784,22 @@
{
/* [84] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [85] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [86] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(107),
+ /* matcher_indices */ MatcherIndicesIndex(110),
},
{
/* [87] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [88] */
@@ -1811,17 +1814,17 @@
{
/* [90] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [91] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(113),
+ /* matcher_indices */ MatcherIndicesIndex(116),
},
{
/* [92] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [93] */
@@ -1836,17 +1839,17 @@
{
/* [95] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [96] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(107),
+ /* matcher_indices */ MatcherIndicesIndex(110),
},
{
/* [97] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [98] */
@@ -1861,17 +1864,17 @@
{
/* [100] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [101] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(113),
+ /* matcher_indices */ MatcherIndicesIndex(116),
},
{
/* [102] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [103] */
@@ -1886,17 +1889,17 @@
{
/* [105] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [106] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(120),
},
{
/* [107] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [108] */
@@ -1916,7 +1919,7 @@
{
/* [111] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(119),
+ /* matcher_indices */ MatcherIndicesIndex(122),
},
{
/* [112] */
@@ -1961,12 +1964,12 @@
{
/* [120] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(120),
},
{
/* [121] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [122] */
@@ -1981,7 +1984,7 @@
{
/* [124] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(119),
+ /* matcher_indices */ MatcherIndicesIndex(122),
},
{
/* [125] */
@@ -2026,7 +2029,7 @@
{
/* [133] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [134] */
@@ -2046,7 +2049,7 @@
{
/* [137] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(123),
+ /* matcher_indices */ MatcherIndicesIndex(126),
},
{
/* [138] */
@@ -2066,7 +2069,7 @@
{
/* [141] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(123),
+ /* matcher_indices */ MatcherIndicesIndex(126),
},
{
/* [142] */
@@ -2081,12 +2084,12 @@
{
/* [144] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(125),
+ /* matcher_indices */ MatcherIndicesIndex(128),
},
{
/* [145] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [146] */
@@ -2101,12 +2104,12 @@
{
/* [148] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(108),
+ /* matcher_indices */ MatcherIndicesIndex(111),
},
{
/* [149] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [150] */
@@ -2121,12 +2124,12 @@
{
/* [152] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(114),
+ /* matcher_indices */ MatcherIndicesIndex(117),
},
{
/* [153] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(123),
+ /* matcher_indices */ MatcherIndicesIndex(126),
},
{
/* [154] */
@@ -2141,12 +2144,12 @@
{
/* [156] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(153),
+ /* matcher_indices */ MatcherIndicesIndex(156),
},
{
/* [157] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [158] */
@@ -2166,7 +2169,7 @@
{
/* [161] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [162] */
@@ -2186,7 +2189,7 @@
{
/* [165] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [166] */
@@ -2206,7 +2209,7 @@
{
/* [169] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [170] */
@@ -2246,7 +2249,7 @@
{
/* [177] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [178] */
@@ -2256,7 +2259,7 @@
{
/* [179] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(111),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [180] */
@@ -2266,7 +2269,7 @@
{
/* [181] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [182] */
@@ -2276,7 +2279,7 @@
{
/* [183] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(111),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [184] */
@@ -2286,7 +2289,7 @@
{
/* [185] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [186] */
@@ -2296,7 +2299,7 @@
{
/* [187] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(135),
+ /* matcher_indices */ MatcherIndicesIndex(138),
},
{
/* [188] */
@@ -2306,7 +2309,7 @@
{
/* [189] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [190] */
@@ -2341,12 +2344,12 @@
{
/* [196] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(107),
+ /* matcher_indices */ MatcherIndicesIndex(110),
},
{
/* [197] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(109),
+ /* matcher_indices */ MatcherIndicesIndex(112),
},
{
/* [198] */
@@ -2356,17 +2359,17 @@
{
/* [199] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [200] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(113),
+ /* matcher_indices */ MatcherIndicesIndex(116),
},
{
/* [201] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [202] */
@@ -2376,17 +2379,17 @@
{
/* [203] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [204] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(117),
+ /* matcher_indices */ MatcherIndicesIndex(120),
},
{
/* [205] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
},
{
/* [206] */
@@ -2401,7 +2404,7 @@
{
/* [208] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(119),
+ /* matcher_indices */ MatcherIndicesIndex(122),
},
{
/* [209] */
@@ -2486,7 +2489,7 @@
{
/* [225] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(139),
+ /* matcher_indices */ MatcherIndicesIndex(142),
},
{
/* [226] */
@@ -2506,7 +2509,7 @@
{
/* [229] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(139),
+ /* matcher_indices */ MatcherIndicesIndex(142),
},
{
/* [230] */
@@ -2526,7 +2529,7 @@
{
/* [233] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(139),
+ /* matcher_indices */ MatcherIndicesIndex(142),
},
{
/* [234] */
@@ -2546,7 +2549,7 @@
{
/* [237] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(141),
+ /* matcher_indices */ MatcherIndicesIndex(144),
},
{
/* [238] */
@@ -2566,7 +2569,7 @@
{
/* [241] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(141),
+ /* matcher_indices */ MatcherIndicesIndex(144),
},
{
/* [242] */
@@ -2586,7 +2589,7 @@
{
/* [245] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(141),
+ /* matcher_indices */ MatcherIndicesIndex(144),
},
{
/* [246] */
@@ -2606,7 +2609,7 @@
{
/* [249] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(141),
+ /* matcher_indices */ MatcherIndicesIndex(144),
},
{
/* [250] */
@@ -2626,7 +2629,7 @@
{
/* [253] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(141),
+ /* matcher_indices */ MatcherIndicesIndex(144),
},
{
/* [254] */
@@ -2646,7 +2649,7 @@
{
/* [257] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(141),
+ /* matcher_indices */ MatcherIndicesIndex(144),
},
{
/* [258] */
@@ -2681,7 +2684,7 @@
{
/* [264] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(111),
+ /* matcher_indices */ MatcherIndicesIndex(114),
},
{
/* [265] */
@@ -2696,7 +2699,7 @@
{
/* [267] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(135),
+ /* matcher_indices */ MatcherIndicesIndex(138),
},
{
/* [268] */
@@ -2711,7 +2714,7 @@
{
/* [270] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(135),
+ /* matcher_indices */ MatcherIndicesIndex(138),
},
{
/* [271] */
@@ -2841,7 +2844,7 @@
{
/* [296] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(108),
+ /* matcher_indices */ MatcherIndicesIndex(111),
},
{
/* [297] */
@@ -2851,7 +2854,7 @@
{
/* [298] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(114),
+ /* matcher_indices */ MatcherIndicesIndex(117),
},
{
/* [299] */
@@ -2861,7 +2864,7 @@
{
/* [300] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(118),
+ /* matcher_indices */ MatcherIndicesIndex(121),
},
{
/* [301] */
@@ -2871,7 +2874,7 @@
{
/* [302] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(120),
+ /* matcher_indices */ MatcherIndicesIndex(123),
},
{
/* [303] */
@@ -2881,12 +2884,12 @@
{
/* [304] */
/* usage */ core::ParameterUsage::kInputAttachment,
- /* matcher_indices */ MatcherIndicesIndex(137),
+ /* matcher_indices */ MatcherIndicesIndex(140),
},
{
/* [305] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(121),
+ /* matcher_indices */ MatcherIndicesIndex(124),
},
{
/* [306] */
@@ -2936,17 +2939,22 @@
{
/* [315] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(143),
+ /* matcher_indices */ MatcherIndicesIndex(146),
},
{
/* [316] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(145),
+ /* matcher_indices */ MatcherIndicesIndex(148),
},
{
/* [317] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(147),
+ /* matcher_indices */ MatcherIndicesIndex(150),
+ },
+ {
+ /* [318] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(107),
},
};
@@ -2957,55 +2965,55 @@
{
/* [0] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(152),
+ /* matcher_indices */ MatcherIndicesIndex(155),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [1] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [2] */
/* name */ "I",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [3] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [4] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(152),
+ /* matcher_indices */ MatcherIndicesIndex(155),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [5] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [6] */
/* name */ "B",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [7] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [8] */
/* name */ "D",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3017,19 +3025,19 @@
{
/* [10] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
/* kind */ TemplateInfo::Kind::kNumber,
},
{
/* [11] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [12] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3041,19 +3049,19 @@
{
/* [14] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
/* kind */ TemplateInfo::Kind::kNumber,
},
{
/* [15] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [16] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3065,25 +3073,25 @@
{
/* [18] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(115),
+ /* matcher_indices */ MatcherIndicesIndex(118),
/* kind */ TemplateInfo::Kind::kNumber,
},
{
/* [19] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [20] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [21] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(151),
+ /* matcher_indices */ MatcherIndicesIndex(154),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3107,7 +3115,7 @@
{
/* [25] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3119,7 +3127,7 @@
{
/* [27] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(150),
+ /* matcher_indices */ MatcherIndicesIndex(153),
/* kind */ TemplateInfo::Kind::kNumber,
},
{
@@ -3131,55 +3139,55 @@
{
/* [29] */
/* name */ "B",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [30] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [31] */
/* name */ "I",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [32] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [33] */
/* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [34] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(152),
+ /* matcher_indices */ MatcherIndicesIndex(155),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [35] */
/* name */ "C",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [36] */
/* name */ "D",
- /* matcher_indices */ MatcherIndicesIndex(149),
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [37] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(151),
+ /* matcher_indices */ MatcherIndicesIndex(154),
/* kind */ TemplateInfo::Kind::kType,
},
{
@@ -3196,15 +3204,15 @@
},
{
/* [40] */
- /* name */ "I",
- /* matcher_indices */ MatcherIndicesIndex(63),
+ /* name */ "R",
+ /* matcher_indices */ MatcherIndicesIndex(152),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [41] */
- /* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(/* invalid */),
- /* kind */ TemplateInfo::Kind::kNumber,
+ /* name */ "T",
+ /* matcher_indices */ MatcherIndicesIndex(152),
+ /* kind */ TemplateInfo::Kind::kType,
},
{
/* [42] */
@@ -3215,14 +3223,14 @@
{
/* [43] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(151),
+ /* matcher_indices */ MatcherIndicesIndex(154),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [44] */
- /* name */ "F",
- /* matcher_indices */ MatcherIndicesIndex(/* invalid */),
- /* kind */ TemplateInfo::Kind::kNumber,
+ /* name */ "I",
+ /* matcher_indices */ MatcherIndicesIndex(63),
+ /* kind */ TemplateInfo::Kind::kType,
},
{
/* [45] */
@@ -3232,26 +3240,38 @@
},
{
/* [46] */
- /* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(152),
- /* kind */ TemplateInfo::Kind::kType,
+ /* name */ "F",
+ /* matcher_indices */ MatcherIndicesIndex(/* invalid */),
+ /* kind */ TemplateInfo::Kind::kNumber,
},
{
/* [47] */
- /* name */ "S",
- /* matcher_indices */ MatcherIndicesIndex(154),
- /* kind */ TemplateInfo::Kind::kType,
+ /* name */ "A",
+ /* matcher_indices */ MatcherIndicesIndex(/* invalid */),
+ /* kind */ TemplateInfo::Kind::kNumber,
},
{
/* [48] */
+ /* name */ "T",
+ /* matcher_indices */ MatcherIndicesIndex(155),
+ /* kind */ TemplateInfo::Kind::kType,
+ },
+ {
+ /* [49] */
+ /* name */ "S",
+ /* matcher_indices */ MatcherIndicesIndex(157),
+ /* kind */ TemplateInfo::Kind::kType,
+ },
+ {
+ /* [50] */
/* name */ "N",
/* matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* kind */ TemplateInfo::Kind::kNumber,
},
{
- /* [49] */
+ /* [51] */
/* name */ "T",
- /* matcher_indices */ MatcherIndicesIndex(155),
+ /* matcher_indices */ MatcherIndicesIndex(158),
/* kind */ TemplateInfo::Kind::kType,
},
};
@@ -3785,7 +3805,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(132),
- /* return_matcher_indices */ MatcherIndicesIndex(131),
+ /* return_matcher_indices */ MatcherIndicesIndex(134),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3796,7 +3816,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(136),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3807,7 +3827,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(140),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3818,7 +3838,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(292),
- /* return_matcher_indices */ MatcherIndicesIndex(131),
+ /* return_matcher_indices */ MatcherIndicesIndex(134),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3829,7 +3849,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(294),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3840,7 +3860,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(0),
/* parameters */ ParameterIndex(144),
- /* return_matcher_indices */ MatcherIndicesIndex(131),
+ /* return_matcher_indices */ MatcherIndicesIndex(134),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3851,7 +3871,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(148),
- /* return_matcher_indices */ MatcherIndicesIndex(131),
+ /* return_matcher_indices */ MatcherIndicesIndex(134),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3862,7 +3882,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(152),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3873,7 +3893,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(300),
- /* return_matcher_indices */ MatcherIndicesIndex(131),
+ /* return_matcher_indices */ MatcherIndicesIndex(134),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3884,7 +3904,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(302),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3895,7 +3915,7 @@
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
/* parameters */ ParameterIndex(156),
- /* return_matcher_indices */ MatcherIndicesIndex(131),
+ /* return_matcher_indices */ MatcherIndicesIndex(134),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3904,7 +3924,7 @@
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(44),
+ /* templates */ TemplateIndex(46),
/* parameters */ ParameterIndex(260),
/* return_matcher_indices */ MatcherIndicesIndex(63),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -3915,9 +3935,9 @@
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(44),
+ /* templates */ TemplateIndex(46),
/* parameters */ ParameterIndex(263),
- /* return_matcher_indices */ MatcherIndicesIndex(131),
+ /* return_matcher_indices */ MatcherIndicesIndex(134),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3926,9 +3946,9 @@
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(44),
+ /* templates */ TemplateIndex(46),
/* parameters */ ParameterIndex(266),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -3937,9 +3957,9 @@
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(44),
+ /* templates */ TemplateIndex(46),
/* parameters */ ParameterIndex(269),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4390,7 +4410,7 @@
/* num_templates */ 2,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(286),
- /* return_matcher_indices */ MatcherIndicesIndex(131),
+ /* return_matcher_indices */ MatcherIndicesIndex(134),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4401,7 +4421,7 @@
/* num_templates */ 2,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(288),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4412,7 +4432,7 @@
/* num_templates */ 2,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(290),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4423,7 +4443,7 @@
/* num_templates */ 2,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(292),
- /* return_matcher_indices */ MatcherIndicesIndex(131),
+ /* return_matcher_indices */ MatcherIndicesIndex(134),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4434,7 +4454,7 @@
/* num_templates */ 2,
/* templates */ TemplateIndex(4),
/* parameters */ ParameterIndex(294),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4445,7 +4465,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
/* parameters */ ParameterIndex(296),
- /* return_matcher_indices */ MatcherIndicesIndex(131),
+ /* return_matcher_indices */ MatcherIndicesIndex(134),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4456,7 +4476,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
/* parameters */ ParameterIndex(298),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4467,7 +4487,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
/* parameters */ ParameterIndex(300),
- /* return_matcher_indices */ MatcherIndicesIndex(131),
+ /* return_matcher_indices */ MatcherIndicesIndex(134),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4478,7 +4498,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
/* parameters */ ParameterIndex(302),
- /* return_matcher_indices */ MatcherIndicesIndex(133),
+ /* return_matcher_indices */ MatcherIndicesIndex(136),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4487,7 +4507,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(46),
+ /* templates */ TemplateIndex(48),
/* parameters */ ParameterIndex(128),
/* return_matcher_indices */ MatcherIndicesIndex(28),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -4498,7 +4518,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(46),
+ /* templates */ TemplateIndex(48),
/* parameters */ ParameterIndex(286),
/* return_matcher_indices */ MatcherIndicesIndex(31),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -4509,7 +4529,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(46),
+ /* templates */ TemplateIndex(48),
/* parameters */ ParameterIndex(288),
/* return_matcher_indices */ MatcherIndicesIndex(34),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -4520,7 +4540,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(46),
+ /* templates */ TemplateIndex(48),
/* parameters */ ParameterIndex(290),
/* return_matcher_indices */ MatcherIndicesIndex(37),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -4531,7 +4551,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(46),
+ /* templates */ TemplateIndex(48),
/* parameters */ ParameterIndex(292),
/* return_matcher_indices */ MatcherIndicesIndex(40),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -4542,7 +4562,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(46),
+ /* templates */ TemplateIndex(48),
/* parameters */ ParameterIndex(294),
/* return_matcher_indices */ MatcherIndicesIndex(43),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -4553,9 +4573,9 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 1,
- /* templates */ TemplateIndex(47),
+ /* templates */ TemplateIndex(49),
/* parameters */ ParameterIndex(296),
- /* return_matcher_indices */ MatcherIndicesIndex(107),
+ /* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4564,9 +4584,9 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 1,
- /* templates */ TemplateIndex(47),
+ /* templates */ TemplateIndex(49),
/* parameters */ ParameterIndex(298),
- /* return_matcher_indices */ MatcherIndicesIndex(113),
+ /* return_matcher_indices */ MatcherIndicesIndex(116),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4575,9 +4595,9 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 1,
- /* templates */ TemplateIndex(47),
+ /* templates */ TemplateIndex(49),
/* parameters */ ParameterIndex(300),
- /* return_matcher_indices */ MatcherIndicesIndex(117),
+ /* return_matcher_indices */ MatcherIndicesIndex(120),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4586,9 +4606,9 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 1,
- /* templates */ TemplateIndex(47),
+ /* templates */ TemplateIndex(49),
/* parameters */ ParameterIndex(302),
- /* return_matcher_indices */ MatcherIndicesIndex(119),
+ /* return_matcher_indices */ MatcherIndicesIndex(122),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4885,7 +4905,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(21),
/* parameters */ ParameterIndex(315),
- /* return_matcher_indices */ MatcherIndicesIndex(143),
+ /* return_matcher_indices */ MatcherIndicesIndex(146),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4896,7 +4916,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(21),
/* parameters */ ParameterIndex(316),
- /* return_matcher_indices */ MatcherIndicesIndex(145),
+ /* return_matcher_indices */ MatcherIndicesIndex(148),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4907,7 +4927,7 @@
/* num_templates */ 1,
/* templates */ TemplateIndex(21),
/* parameters */ ParameterIndex(317),
- /* return_matcher_indices */ MatcherIndicesIndex(147),
+ /* return_matcher_indices */ MatcherIndicesIndex(150),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -4916,7 +4936,7 @@
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
/* num_templates */ 1,
- /* templates */ TemplateIndex(49),
+ /* templates */ TemplateIndex(51),
/* parameters */ ParameterIndex(275),
/* return_matcher_indices */ MatcherIndicesIndex(3),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -4927,7 +4947,7 @@
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(48),
+ /* templates */ TemplateIndex(50),
/* parameters */ ParameterIndex(278),
/* return_matcher_indices */ MatcherIndicesIndex(25),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -4957,16 +4977,38 @@
{
/* [154] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 1,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(40),
+ /* parameters */ ParameterIndex(1),
+ /* return_matcher_indices */ MatcherIndicesIndex(3),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [155] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 1,
+ /* num_templates */ 3,
+ /* templates */ TemplateIndex(40),
+ /* parameters */ ParameterIndex(318),
+ /* return_matcher_indices */ MatcherIndicesIndex(98),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [156] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(40),
+ /* templates */ TemplateIndex(44),
/* parameters */ ParameterIndex(284),
/* return_matcher_indices */ MatcherIndicesIndex(63),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [155] */
+ /* [157] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -4977,7 +5019,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [156] */
+ /* [158] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 6,
/* num_explicit_templates */ 0,
@@ -4988,7 +5030,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [157] */
+ /* [159] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -4999,7 +5041,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [158] */
+ /* [160] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -5010,7 +5052,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [159] */
+ /* [161] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5021,7 +5063,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [160] */
+ /* [162] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5032,7 +5074,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [161] */
+ /* [163] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5043,7 +5085,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [162] */
+ /* [164] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5054,7 +5096,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [163] */
+ /* [165] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5065,7 +5107,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [164] */
+ /* [166] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5076,7 +5118,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [165] */
+ /* [167] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -5087,7 +5129,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [166] */
+ /* [168] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -5107,91 +5149,91 @@
/* [0] */
/* fn array_length[I : u32, A : access](ptr<storage, struct_with_runtime_array, A>, I) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(154),
+ /* overloads */ OverloadIndex(156),
},
{
/* [1] */
/* fn atomic_and[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(155),
+ /* overloads */ OverloadIndex(157),
},
{
/* [2] */
/* fn atomic_compare_exchange[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, U, T, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(156),
+ /* overloads */ OverloadIndex(158),
},
{
/* [3] */
/* fn atomic_exchange[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(155),
+ /* overloads */ OverloadIndex(157),
},
{
/* [4] */
/* fn atomic_iadd[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(155),
+ /* overloads */ OverloadIndex(157),
},
{
/* [5] */
/* fn atomic_isub[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(155),
+ /* overloads */ OverloadIndex(157),
},
{
/* [6] */
/* fn atomic_load[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(157),
+ /* overloads */ OverloadIndex(159),
},
{
/* [7] */
/* fn atomic_or[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(155),
+ /* overloads */ OverloadIndex(157),
},
{
/* [8] */
/* fn atomic_smax[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(155),
+ /* overloads */ OverloadIndex(157),
},
{
/* [9] */
/* fn atomic_smin[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(155),
+ /* overloads */ OverloadIndex(157),
},
{
/* [10] */
/* fn atomic_store[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(158),
+ /* overloads */ OverloadIndex(160),
},
{
/* [11] */
/* fn atomic_umax[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(155),
+ /* overloads */ OverloadIndex(157),
},
{
/* [12] */
/* fn atomic_umin[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(155),
+ /* overloads */ OverloadIndex(157),
},
{
/* [13] */
/* fn atomic_xor[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, U, U, T) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(155),
+ /* overloads */ OverloadIndex(157),
},
{
/* [14] */
/* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(159),
+ /* overloads */ OverloadIndex(161),
},
{
/* [15] */
@@ -5389,19 +5431,19 @@
/* [26] */
/* fn matrix_times_matrix[T : f32_f16, K : num, C : num, R : num](mat<K, R, T>, mat<C, K, T>) -> mat<C, R, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(160),
+ /* overloads */ OverloadIndex(162),
},
{
/* [27] */
/* fn matrix_times_scalar[T : f32_f16, N : num, M : num](mat<N, M, T>, T) -> mat<N, M, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(161),
+ /* overloads */ OverloadIndex(163),
},
{
/* [28] */
/* fn matrix_times_vector[T : f32_f16, N : num, M : num](mat<N, M, T>, vec<N, T>) -> vec<M, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(162),
+ /* overloads */ OverloadIndex(164),
},
{
/* [29] */
@@ -5429,13 +5471,13 @@
/* [31] */
/* fn vector_times_matrix[T : f32_f16, N : num, M : num](vec<N, T>, mat<M, N, T>) -> vec<M, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(163),
+ /* overloads */ OverloadIndex(165),
},
{
/* [32] */
/* fn vector_times_scalar[T : f32_f16, N : num](vec<N, T>, T) -> vec<N, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(164),
+ /* overloads */ OverloadIndex(166),
},
{
/* [33] */
@@ -5454,15 +5496,22 @@
},
{
/* [35] */
- /* fn sdot(u32, u32, u32) -> i32 */
- /* num overloads */ 1,
- /* overloads */ OverloadIndex(165),
+ /* fn sign<R : iu32>[T : iu32](T) -> R */
+ /* fn sign<R : iu32>[T : iu32, N : num](vec<N, T>) -> vec<N, R> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(154),
},
{
/* [36] */
+ /* fn sdot(u32, u32, u32) -> i32 */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(167),
+ },
+ {
+ /* [37] */
/* fn udot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(166),
+ /* overloads */ OverloadIndex(168),
},
};
diff --git a/src/tint/lang/spirv/ir/builtin_call.cc b/src/tint/lang/spirv/ir/builtin_call.cc
index ea2fd2a..8960fed 100644
--- a/src/tint/lang/spirv/ir/builtin_call.cc
+++ b/src/tint/lang/spirv/ir/builtin_call.cc
@@ -27,8 +27,6 @@
#include "src/tint/lang/spirv/ir/builtin_call.h"
-#include <utility>
-
#include "src/tint/lang/core/ir/clone_context.h"
#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/ice/ice.h"
@@ -51,11 +49,21 @@
BuiltinCall* BuiltinCall::Clone(core::ir::CloneContext& ctx) {
auto* new_result = ctx.Clone(Result(0));
auto new_args = ctx.Clone<BuiltinCall::kDefaultNumOperands>(Args());
- return ctx.ir.CreateInstruction<BuiltinCall>(new_result, func_, new_args);
+ auto new_explicit = ExplicitTemplateParams();
+
+ auto* inst = ctx.ir.CreateInstruction<BuiltinCall>(new_result, func_, new_args);
+ if (!new_explicit.IsEmpty()) {
+ inst->SetExplicitTemplateParams(new_explicit);
+ }
+ return inst;
}
tint::core::ir::Instruction::Accesses BuiltinCall::GetSideEffects() const {
return spirv::GetSideEffects(func_);
}
+std::string BuiltinCall::FriendlyName() const {
+ return std::string("spirv.") + str(func_) + ExplicitTemplateParamsToStr();
+}
+
} // namespace tint::spirv::ir
diff --git a/src/tint/lang/spirv/ir/builtin_call.h b/src/tint/lang/spirv/ir/builtin_call.h
index 733b3da..8d9e44f 100644
--- a/src/tint/lang/spirv/ir/builtin_call.h
+++ b/src/tint/lang/spirv/ir/builtin_call.h
@@ -63,7 +63,7 @@
size_t FuncId() const override { return static_cast<size_t>(func_); }
/// @returns the friendly name for the instruction
- std::string FriendlyName() const override { return std::string("spirv.") + str(func_); }
+ std::string FriendlyName() const override;
/// @returns the table data to validate this builtin
const core::intrinsic::TableData& TableData() const override {
diff --git a/src/tint/lang/spirv/ir/builtin_call_test.cc b/src/tint/lang/spirv/ir/builtin_call_test.cc
index bd5ba07..211aa30 100644
--- a/src/tint/lang/spirv/ir/builtin_call_test.cc
+++ b/src/tint/lang/spirv/ir/builtin_call_test.cc
@@ -46,6 +46,7 @@
EXPECT_EQ(mod.Types().f32(), new_b->Result(0)->Type());
EXPECT_EQ(BuiltinFn::kArrayLength, new_b->Func());
+ EXPECT_TRUE(new_b->ExplicitTemplateParams().IsEmpty());
auto args = new_b->Args();
EXPECT_EQ(2u, args.Length());
@@ -57,6 +58,35 @@
EXPECT_EQ(2_u, val1->As<core::constant::Scalar<core::u32>>()->ValueAs<core::u32>());
}
+TEST_F(IR_SpirvBuiltinCallTest, CloneWithExplicitParams) {
+ auto* builtin = b.Call<BuiltinCall>(mod.Types().f32(), BuiltinFn::kArrayLength, 1_u, 2_u);
+ builtin->SetExplicitTemplateParams(
+ Vector<const core::type::Type*, 2>{mod.Types().f32(), mod.Types().i32()});
+
+ auto* new_b = clone_ctx.Clone(builtin);
+
+ EXPECT_NE(builtin, new_b);
+ EXPECT_NE(builtin->Result(0), new_b->Result(0));
+ EXPECT_EQ(mod.Types().f32(), new_b->Result(0)->Type());
+
+ EXPECT_EQ(BuiltinFn::kArrayLength, new_b->Func());
+
+ auto args = new_b->Args();
+ EXPECT_EQ(2u, args.Length());
+
+ auto* val0 = args[0]->As<core::ir::Constant>()->Value();
+ EXPECT_EQ(1_u, val0->As<core::constant::Scalar<core::u32>>()->ValueAs<core::u32>());
+
+ auto* val1 = args[1]->As<core::ir::Constant>()->Value();
+ EXPECT_EQ(2_u, val1->As<core::constant::Scalar<core::u32>>()->ValueAs<core::u32>());
+
+ auto new_explicit = new_b->ExplicitTemplateParams();
+ ASSERT_EQ(2u, new_explicit.Length());
+
+ EXPECT_EQ(mod.Types().f32(), new_explicit[0]);
+ EXPECT_EQ(mod.Types().i32(), new_explicit[1]);
+}
+
TEST_F(IR_SpirvBuiltinCallTest, CloneNoArgs) {
auto* builtin = b.Call<BuiltinCall>(mod.Types().f32(), BuiltinFn::kArrayLength);
@@ -65,6 +95,7 @@
EXPECT_EQ(mod.Types().f32(), new_b->Result(0)->Type());
EXPECT_EQ(BuiltinFn::kArrayLength, new_b->Func());
+ EXPECT_TRUE(new_b->ExplicitTemplateParams().IsEmpty());
auto args = new_b->Args();
EXPECT_TRUE(args.IsEmpty());
diff --git a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
index da30cd1..f2cc8f3 100644
--- a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
+++ b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
@@ -502,7 +502,7 @@
)");
}
-TEST_P(SpirvReaderTest_GlslStd450_Inting_Inting_SignednessCoercing, DISABLED_Scalar_UnsignedArg) {
+TEST_P(SpirvReaderTest_GlslStd450_Inting_Inting_SignednessCoercing, Scalar_UnsignedArg) {
EXPECT_IR(Preamble() + R"(
%1 = OpExtInst %int %glsl )" +
GetParam().opcode +
@@ -514,27 +514,31 @@
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- let x_1 = )" + GetParam().wgsl_func +
- R"((bitcast<i32>(u1));
+ %2:i32 = bitcast 10u
+ %3:i32 = sign %2
+ %4:i32 = let %3
+ ret
}
}
)");
}
-TEST_P(SpirvReaderTest_GlslStd450_Inting_Inting_SignednessCoercing,
- DISABLED_Scalar_UnsignedResult) {
+TEST_P(SpirvReaderTest_GlslStd450_Inting_Inting_SignednessCoercing, Scalar_UnsignedResult) {
EXPECT_IR(Preamble() + R"(
%1 = OpExtInst %uint %glsl )" +
GetParam().opcode +
R"( %int_30
+ %2 = OpCopyObject %uint %1
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- let x_1 = bitcast<u32>()" +
- GetParam().wgsl_func + R"((i1));
+ %2:i32 = sign 30i
+ %3:u32 = bitcast %2
+ %4:u32 = let %3
+ ret
}
}
)");
@@ -562,38 +566,43 @@
)");
}
-TEST_P(SpirvReaderTest_GlslStd450_Inting_Inting_SignednessCoercing, DISABLED_Vector_UnsignedArg) {
+TEST_P(SpirvReaderTest_GlslStd450_Inting_Inting_SignednessCoercing, Vector_UnsignedArg) {
EXPECT_IR(Preamble() + R"(
%1 = OpExtInst %v2int %glsl )" +
GetParam().opcode +
R"( %v2uint_10_20
+ %2 = OpCopyObject %v2int %1
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- let x_1 = )" + GetParam().wgsl_func +
- R"((bitcast<vec2i>(v2u1));
+ %2:vec2<i32> = bitcast vec2<u32>(10u, 20u)
+ %3:vec2<i32> = sign %2
+ %4:vec2<i32> = let %3
+ ret
}
}
)");
}
-TEST_P(SpirvReaderTest_GlslStd450_Inting_Inting_SignednessCoercing,
- DISABLED_Vector_UnsignedResult) {
+TEST_P(SpirvReaderTest_GlslStd450_Inting_Inting_SignednessCoercing, Vector_UnsignedResult) {
EXPECT_IR(Preamble() + R"(
%1 = OpExtInst %v2uint %glsl )" +
GetParam().opcode +
R"( %v2int_30_40
+ %2 = OpCopyObject %v2uint %1
OpReturn
OpFunctionEnd
)",
R"(
%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
$B1: {
- let x_1 = bitcast<vec2u>()" +
- GetParam().wgsl_func + R"((v2i1));
+ %2:vec2<i32> = sign vec2<i32>(30i, 40i)
+ %3:vec2<u32> = bitcast %2
+ %4:vec2<u32> = let %3
+ ret
}
}
)");
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index b59b33a..4746914 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -67,12 +67,40 @@
case spirv::BuiltinFn::kInverse:
Inverse(builtin);
break;
+ case spirv::BuiltinFn::kSign:
+ Sign(builtin);
+ break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
}
}
}
+ // The `spirv.sign` method takes a `u32` operand which is not accepted in WGSL. If the operand
+ // is a `u32` or a `vec<N, u32>` then we need to bitcast the operand to `i32` before doing the
+ // comparison.
+ void Sign(spirv::ir::BuiltinCall* call) {
+ auto* arg = call->Args()[0];
+
+ b.InsertBefore(call, [&] {
+ auto* result_ty = call->Result(0)->Type();
+ if (arg->Type()->IsUnsignedIntegerScalarOrVector()) {
+ arg = b.Bitcast(ty.MatchWidth(ty.i32(), result_ty), arg)->Result(0);
+ }
+ auto* new_call =
+ b.Call(result_ty, core::BuiltinFn::kSign, Vector<core::ir::Value*, 1>{arg});
+
+ core::ir::Value* replacement = new_call->Result(0);
+ // If the call is a `u32` result type, we need to cast it to `i32`.
+ if (result_ty->DeepestElement() == ty.u32()) {
+ new_call->Result(0)->SetType(ty.MatchWidth(ty.i32(), result_ty));
+ replacement = b.Bitcast(result_ty, replacement)->Result(0);
+ }
+ call->Result(0)->ReplaceAllUsesWith(replacement);
+ });
+ call->Destroy();
+ }
+
void Normalize(spirv::ir::BuiltinCall* call) {
auto* arg = call->Args()[0];
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index 7f0d1c7..d622215 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -739,5 +739,269 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvParser_BuiltinsTest, SSign_Scalar_UnsignedArg) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kSign,
+ Vector<const core::type::Type*, 1>{ty.i32()},
+ b.Constant(10_u));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.sign<i32> 10u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = bitcast 10u
+ %3:i32 = sign %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, SSign_Scalar_UnsignedResult) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kSign,
+ Vector<const core::type::Type*, 1>{ty.u32()},
+ b.Constant(10_i));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.sign<u32> 10i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = sign 10i
+ %3:u32 = bitcast %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, SSign_Scalar_UnsignedArgAndResult) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kSign,
+ Vector<const core::type::Type*, 1>{ty.u32()},
+ b.Constant(10_u));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.sign<u32> 10u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = bitcast 10u
+ %3:i32 = sign %2
+ %4:u32 = bitcast %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, SSign_Scalar_SignedArgAndResult) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kSign,
+ Vector<const core::type::Type*, 1>{ty.i32()},
+ b.Constant(10_i));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.sign<i32> 10i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = sign 10i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, SSign_Vector_UnsignedArg) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kSign,
+ Vector<const core::type::Type*, 1>{ty.i32()},
+ b.Splat(ty.vec2<u32>(), (10_u)));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.sign<i32> vec2<u32>(10u)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = bitcast vec2<u32>(10u)
+ %3:vec2<i32> = sign %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, SSign_Vector_UnsignedResult) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kSign,
+ Vector<const core::type::Type*, 1>{ty.u32()},
+ b.Splat(ty.vec2<i32>(), 10_i));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.sign<u32> vec2<i32>(10i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = sign vec2<i32>(10i)
+ %3:vec2<u32> = bitcast %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, SSign_Vector_UnsignedArgAndResult) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kSign,
+ Vector<const core::type::Type*, 1>{ty.u32()},
+ b.Splat(ty.vec2<u32>(), 10_u));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.sign<u32> vec2<u32>(10u)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = bitcast vec2<u32>(10u)
+ %3:vec2<i32> = sign %2
+ %4:vec2<u32> = bitcast %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, SSign_Vector_SignedArgAndResult) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kSign,
+ Vector<const core::type::Type*, 1>{ty.i32()},
+ b.Splat(ty.vec2<i32>(), 10_i));
+ b.Return(ep);
+ });
+
+ auto* src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.sign<i32> vec2<i32>(10i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto* expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = sign vec2<i32>(10i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index fe50b39..db2d433 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -667,7 +667,6 @@
case GLSLstd450SAbs:
return core::BuiltinFn::kAbs;
case GLSLstd450FSign:
- case GLSLstd450SSign:
return core::BuiltinFn::kSign;
case GLSLstd450FindILsb:
return core::BuiltinFn::kFirstTrailingBit;
@@ -759,6 +758,8 @@
spirv::BuiltinFn GetGlslStd450SpirvEquivalentFuncName(uint32_t ext_opcode) {
switch (ext_opcode) {
+ case GLSLstd450SSign:
+ return spirv::BuiltinFn::kSign;
case GLSLstd450Normalize:
return spirv::BuiltinFn::kNormalize;
case GLSLstd450MatrixInverse:
@@ -769,6 +770,14 @@
return spirv::BuiltinFn::kNone;
}
+ Vector<const core::type::Type*, 1> GlslStd450ExplicitParams(uint32_t ext_opcode,
+ const core::type::Type* result_ty) {
+ if (ext_opcode != GLSLstd450SSign) {
+ return {};
+ }
+ return {result_ty->DeepestElement()};
+ }
+
/// @param inst the SPIR-V instruction for OpAccessChain
void EmitGlslStd450ExtInst(const spvtools::opt::Instruction& inst) {
const auto ext_opcode = inst.GetSingleWordInOperand(1);
@@ -788,7 +797,10 @@
const auto spv_fn = GetGlslStd450SpirvEquivalentFuncName(ext_opcode);
if (spv_fn != spirv::BuiltinFn::kNone) {
- Emit(b_.Call<spirv::ir::BuiltinCall>(result_ty, spv_fn, operands), inst.result_id());
+ auto explicit_params = GlslStd450ExplicitParams(ext_opcode, result_ty);
+ Emit(b_.CallExplicit<spirv::ir::BuiltinCall>(result_ty, spv_fn, explicit_params,
+ operands),
+ inst.result_id());
return;
}
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index db03a8d..39aeff1 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -325,6 +325,9 @@
implicit(T: f32_f16) fn inverse(mat3x3<T>) -> mat3x3<T>
implicit(T: f32_f16) fn inverse(mat4x4<T>) -> mat4x4<T>
+implicit(T: iu32) fn sign<R: iu32>(T) -> R
+implicit(T: iu32, N: num) fn sign<R: iu32>(vec<N, T>) -> vec<N, R>
+
////////////////////////////////////////////////////////////////////////////////
// SPV_KHR_integer_dot_product instructions
////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index a25edfc..a37bccd 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1390,6 +1390,9 @@
case spirv::BuiltinFn::kSelect:
op = spv::Op::OpSelect;
break;
+ case spirv::BuiltinFn::kSign:
+ ext_inst(GLSLstd450SSign);
+ break;
case spirv::BuiltinFn::kInverse:
ext_inst(GLSLstd450MatrixInverse);
break;
