[spirv-reader][ir] Correctly handle GLSL 450 FaceForward
The SPIR-V `FaceForward` method allows a scalar values. The WGSL version
requires vectors. In the case of scalar values, do the calculation
manually.
Bug: 391488194
Change-Id: Ib58f809c867e0ba9703fabab524c03b0f434bc59
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/223114
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index 97e23b1..43db555 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -136,6 +136,8 @@
return "findUMsb";
case BuiltinFn::kRefract:
return "refract";
+ case BuiltinFn::kFaceForward:
+ return "faceForward";
case BuiltinFn::kSdot:
return "sdot";
case BuiltinFn::kUdot:
@@ -203,6 +205,7 @@
case BuiltinFn::kFindSMsb:
case BuiltinFn::kFindUMsb:
case BuiltinFn::kRefract:
+ case BuiltinFn::kFaceForward:
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 7062ea2..03ba827 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -86,6 +86,7 @@
case BuiltinFn::kFindSMsb:
case BuiltinFn::kFindUMsb:
case BuiltinFn::kRefract:
+ case BuiltinFn::kFaceForward:
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 099c314..2f4a4e1 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -95,6 +95,7 @@
kFindSMsb,
kFindUMsb,
kRefract,
+ kFaceForward,
kSdot,
kUdot,
kNone,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 5307ca9..8af4dce 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -2837,52 +2837,52 @@
{
/* [290] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(63),
+ /* matcher_indices */ MatcherIndicesIndex(101),
},
{
/* [291] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(63),
+ /* matcher_indices */ MatcherIndicesIndex(101),
},
{
/* [292] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(63),
+ /* matcher_indices */ MatcherIndicesIndex(101),
},
{
/* [293] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(5),
+ /* matcher_indices */ MatcherIndicesIndex(17),
},
{
/* [294] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(3),
+ /* matcher_indices */ MatcherIndicesIndex(63),
},
{
/* [295] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(32),
+ /* matcher_indices */ MatcherIndicesIndex(63),
},
{
/* [296] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(8),
+ /* matcher_indices */ MatcherIndicesIndex(63),
},
{
/* [297] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(35),
+ /* matcher_indices */ MatcherIndicesIndex(5),
},
{
/* [298] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(8),
+ /* matcher_indices */ MatcherIndicesIndex(3),
},
{
/* [299] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(38),
+ /* matcher_indices */ MatcherIndicesIndex(32),
},
{
/* [300] */
@@ -2892,7 +2892,7 @@
{
/* [301] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(41),
+ /* matcher_indices */ MatcherIndicesIndex(35),
},
{
/* [302] */
@@ -2902,7 +2902,7 @@
{
/* [303] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(44),
+ /* matcher_indices */ MatcherIndicesIndex(38),
},
{
/* [304] */
@@ -2912,27 +2912,27 @@
{
/* [305] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(132),
+ /* matcher_indices */ MatcherIndicesIndex(41),
},
{
/* [306] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(3),
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [307] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(138),
+ /* matcher_indices */ MatcherIndicesIndex(44),
},
{
/* [308] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(3),
+ /* matcher_indices */ MatcherIndicesIndex(8),
},
{
/* [309] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(142),
+ /* matcher_indices */ MatcherIndicesIndex(132),
},
{
/* [310] */
@@ -2942,7 +2942,7 @@
{
/* [311] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(144),
+ /* matcher_indices */ MatcherIndicesIndex(138),
},
{
/* [312] */
@@ -2951,77 +2951,92 @@
},
{
/* [313] */
- /* usage */ core::ParameterUsage::kInputAttachment,
- /* matcher_indices */ MatcherIndicesIndex(161),
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(142),
},
{
/* [314] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(145),
+ /* matcher_indices */ MatcherIndicesIndex(3),
},
{
/* [315] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(13),
+ /* matcher_indices */ MatcherIndicesIndex(144),
},
{
/* [316] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(17),
+ /* matcher_indices */ MatcherIndicesIndex(3),
},
{
/* [317] */
- /* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(21),
+ /* usage */ core::ParameterUsage::kInputAttachment,
+ /* matcher_indices */ MatcherIndicesIndex(161),
},
{
/* [318] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(3),
+ /* matcher_indices */ MatcherIndicesIndex(145),
},
{
/* [319] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(21),
+ /* matcher_indices */ MatcherIndicesIndex(13),
},
{
/* [320] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(101),
+ /* matcher_indices */ MatcherIndicesIndex(17),
},
{
/* [321] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(17),
+ /* matcher_indices */ MatcherIndicesIndex(21),
},
{
/* [322] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(113),
+ /* matcher_indices */ MatcherIndicesIndex(3),
},
{
/* [323] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(116),
+ /* matcher_indices */ MatcherIndicesIndex(21),
},
{
/* [324] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(167),
+ /* matcher_indices */ MatcherIndicesIndex(101),
},
{
/* [325] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(169),
+ /* matcher_indices */ MatcherIndicesIndex(113),
},
{
/* [326] */
/* usage */ core::ParameterUsage::kNone,
- /* matcher_indices */ MatcherIndicesIndex(171),
+ /* matcher_indices */ MatcherIndicesIndex(116),
},
{
/* [327] */
/* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(167),
+ },
+ {
+ /* [328] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(169),
+ },
+ {
+ /* [329] */
+ /* usage */ core::ParameterUsage::kNone,
+ /* matcher_indices */ MatcherIndicesIndex(171),
+ },
+ {
+ /* [330] */
+ /* usage */ core::ParameterUsage::kNone,
/* matcher_indices */ MatcherIndicesIndex(107),
},
};
@@ -3947,7 +3962,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(9),
- /* parameters */ ParameterIndex(301),
+ /* parameters */ ParameterIndex(305),
/* return_matcher_indices */ MatcherIndicesIndex(155),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -3958,7 +3973,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(9),
- /* parameters */ ParameterIndex(303),
+ /* parameters */ ParameterIndex(307),
/* return_matcher_indices */ MatcherIndicesIndex(157),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4002,7 +4017,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(309),
+ /* parameters */ ParameterIndex(313),
/* return_matcher_indices */ MatcherIndicesIndex(155),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4013,7 +4028,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(311),
+ /* parameters */ ParameterIndex(315),
/* return_matcher_indices */ MatcherIndicesIndex(157),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4365,7 +4380,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(9),
- /* parameters */ ParameterIndex(313),
+ /* parameters */ ParameterIndex(317),
/* return_matcher_indices */ MatcherIndicesIndex(64),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4519,7 +4534,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(13),
- /* parameters */ ParameterIndex(295),
+ /* parameters */ ParameterIndex(299),
/* return_matcher_indices */ MatcherIndicesIndex(155),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4530,7 +4545,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(13),
- /* parameters */ ParameterIndex(297),
+ /* parameters */ ParameterIndex(301),
/* return_matcher_indices */ MatcherIndicesIndex(157),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4541,7 +4556,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(13),
- /* parameters */ ParameterIndex(299),
+ /* parameters */ ParameterIndex(303),
/* return_matcher_indices */ MatcherIndicesIndex(157),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4552,7 +4567,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(13),
- /* parameters */ ParameterIndex(301),
+ /* parameters */ ParameterIndex(305),
/* return_matcher_indices */ MatcherIndicesIndex(155),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4563,7 +4578,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(13),
- /* parameters */ ParameterIndex(303),
+ /* parameters */ ParameterIndex(307),
/* return_matcher_indices */ MatcherIndicesIndex(157),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4574,7 +4589,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(14),
- /* parameters */ ParameterIndex(305),
+ /* parameters */ ParameterIndex(309),
/* return_matcher_indices */ MatcherIndicesIndex(155),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4585,7 +4600,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(14),
- /* parameters */ ParameterIndex(307),
+ /* parameters */ ParameterIndex(311),
/* return_matcher_indices */ MatcherIndicesIndex(157),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4596,7 +4611,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(14),
- /* parameters */ ParameterIndex(309),
+ /* parameters */ ParameterIndex(313),
/* return_matcher_indices */ MatcherIndicesIndex(155),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4607,7 +4622,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(14),
- /* parameters */ ParameterIndex(311),
+ /* parameters */ ParameterIndex(315),
/* return_matcher_indices */ MatcherIndicesIndex(157),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4629,7 +4644,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(57),
- /* parameters */ ParameterIndex(295),
+ /* parameters */ ParameterIndex(299),
/* return_matcher_indices */ MatcherIndicesIndex(31),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4640,7 +4655,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(57),
- /* parameters */ ParameterIndex(297),
+ /* parameters */ ParameterIndex(301),
/* return_matcher_indices */ MatcherIndicesIndex(34),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4651,7 +4666,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(57),
- /* parameters */ ParameterIndex(299),
+ /* parameters */ ParameterIndex(303),
/* return_matcher_indices */ MatcherIndicesIndex(37),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4662,7 +4677,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(57),
- /* parameters */ ParameterIndex(301),
+ /* parameters */ ParameterIndex(305),
/* return_matcher_indices */ MatcherIndicesIndex(40),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4673,7 +4688,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(57),
- /* parameters */ ParameterIndex(303),
+ /* parameters */ ParameterIndex(307),
/* return_matcher_indices */ MatcherIndicesIndex(43),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4684,7 +4699,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(58),
- /* parameters */ ParameterIndex(305),
+ /* parameters */ ParameterIndex(309),
/* return_matcher_indices */ MatcherIndicesIndex(131),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4695,7 +4710,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(58),
- /* parameters */ ParameterIndex(307),
+ /* parameters */ ParameterIndex(311),
/* return_matcher_indices */ MatcherIndicesIndex(137),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4706,7 +4721,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(58),
- /* parameters */ ParameterIndex(309),
+ /* parameters */ ParameterIndex(313),
/* return_matcher_indices */ MatcherIndicesIndex(141),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -4717,7 +4732,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(58),
- /* parameters */ ParameterIndex(311),
+ /* parameters */ ParameterIndex(315),
/* return_matcher_indices */ MatcherIndicesIndex(143),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5014,7 +5029,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(324),
+ /* parameters */ ParameterIndex(327),
/* return_matcher_indices */ MatcherIndicesIndex(167),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5025,7 +5040,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(325),
+ /* parameters */ ParameterIndex(328),
/* return_matcher_indices */ MatcherIndicesIndex(169),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5036,7 +5051,7 @@
/* num_explicit_templates */ 0,
/* num_templates */ 1,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(326),
+ /* parameters */ ParameterIndex(329),
/* return_matcher_indices */ MatcherIndicesIndex(171),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5102,7 +5117,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 3,
/* templates */ TemplateIndex(50),
- /* parameters */ ParameterIndex(327),
+ /* parameters */ ParameterIndex(330),
/* return_matcher_indices */ MatcherIndicesIndex(98),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5124,7 +5139,7 @@
/* num_explicit_templates */ 1,
/* num_templates */ 4,
/* templates */ TemplateIndex(30),
- /* parameters */ ParameterIndex(322),
+ /* parameters */ ParameterIndex(325),
/* return_matcher_indices */ MatcherIndicesIndex(110),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5175,16 +5190,38 @@
{
/* [162] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 2,
+ /* num_parameters */ 3,
/* num_explicit_templates */ 0,
- /* num_templates */ 2,
- /* templates */ TemplateIndex(53),
- /* parameters */ ParameterIndex(293),
- /* return_matcher_indices */ MatcherIndicesIndex(63),
+ /* num_templates */ 1,
+ /* templates */ TemplateIndex(5),
+ /* parameters */ ParameterIndex(4),
+ /* return_matcher_indices */ MatcherIndicesIndex(3),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
/* [163] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 3,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(47),
+ /* parameters */ ParameterIndex(290),
+ /* return_matcher_indices */ MatcherIndicesIndex(101),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [164] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 2,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(53),
+ /* parameters */ ParameterIndex(297),
+ /* return_matcher_indices */ MatcherIndicesIndex(63),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [165] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -5195,7 +5232,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [164] */
+ /* [166] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 6,
/* num_explicit_templates */ 0,
@@ -5206,7 +5243,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [165] */
+ /* [167] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -5217,7 +5254,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [166] */
+ /* [168] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -5228,7 +5265,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [167] */
+ /* [169] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -5239,36 +5276,14 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [168] */
+ /* [170] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 4,
/* templates */ TemplateIndex(5),
- /* parameters */ ParameterIndex(315),
- /* return_matcher_indices */ MatcherIndicesIndex(9),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [169] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 2,
- /* num_explicit_templates */ 0,
- /* num_templates */ 3,
- /* templates */ TemplateIndex(47),
- /* parameters */ ParameterIndex(317),
- /* return_matcher_indices */ MatcherIndicesIndex(21),
- /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
- },
- {
- /* [170] */
- /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 2,
- /* num_explicit_templates */ 0,
- /* num_templates */ 3,
- /* templates */ TemplateIndex(47),
/* parameters */ ParameterIndex(319),
- /* return_matcher_indices */ MatcherIndicesIndex(98),
+ /* return_matcher_indices */ MatcherIndicesIndex(9),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -5278,8 +5293,8 @@
/* num_explicit_templates */ 0,
/* num_templates */ 3,
/* templates */ TemplateIndex(47),
- /* parameters */ ParameterIndex(320),
- /* return_matcher_indices */ MatcherIndicesIndex(98),
+ /* parameters */ ParameterIndex(321),
+ /* return_matcher_indices */ MatcherIndicesIndex(21),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -5287,6 +5302,28 @@
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
+ /* num_templates */ 3,
+ /* templates */ TemplateIndex(47),
+ /* parameters */ ParameterIndex(323),
+ /* return_matcher_indices */ MatcherIndicesIndex(98),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [173] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 2,
+ /* num_explicit_templates */ 0,
+ /* num_templates */ 3,
+ /* templates */ TemplateIndex(47),
+ /* parameters */ ParameterIndex(292),
+ /* return_matcher_indices */ MatcherIndicesIndex(98),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [174] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 2,
+ /* num_explicit_templates */ 0,
/* num_templates */ 2,
/* templates */ TemplateIndex(47),
/* parameters */ ParameterIndex(288),
@@ -5294,24 +5331,24 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [173] */
+ /* [175] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(290),
+ /* parameters */ ParameterIndex(294),
/* return_matcher_indices */ MatcherIndicesIndex(59),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [174] */
+ /* [176] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
/* num_templates */ 0,
/* templates */ TemplateIndex(/* invalid */),
- /* parameters */ ParameterIndex(290),
+ /* parameters */ ParameterIndex(294),
/* return_matcher_indices */ MatcherIndicesIndex(63),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
@@ -5325,91 +5362,91 @@
/* [0] */
/* fn array_length[I : u32, A : access](ptr<storage, struct_with_runtime_array, A>, I) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(162),
+ /* overloads */ OverloadIndex(164),
},
{
/* [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(163),
+ /* overloads */ OverloadIndex(165),
},
{
/* [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(164),
+ /* overloads */ OverloadIndex(166),
},
{
/* [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(163),
+ /* overloads */ OverloadIndex(165),
},
{
/* [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(163),
+ /* overloads */ OverloadIndex(165),
},
{
/* [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(163),
+ /* overloads */ OverloadIndex(165),
},
{
/* [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(165),
+ /* overloads */ OverloadIndex(167),
},
{
/* [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(163),
+ /* overloads */ OverloadIndex(165),
},
{
/* [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(163),
+ /* overloads */ OverloadIndex(165),
},
{
/* [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(163),
+ /* overloads */ OverloadIndex(165),
},
{
/* [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(166),
+ /* overloads */ OverloadIndex(168),
},
{
/* [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(163),
+ /* overloads */ OverloadIndex(165),
},
{
/* [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(163),
+ /* overloads */ OverloadIndex(165),
},
{
/* [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(163),
+ /* overloads */ OverloadIndex(165),
},
{
/* [14] */
/* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(167),
+ /* overloads */ OverloadIndex(169),
},
{
/* [15] */
@@ -5607,19 +5644,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(168),
+ /* overloads */ OverloadIndex(170),
},
{
/* [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(169),
+ /* overloads */ OverloadIndex(171),
},
{
/* [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(170),
+ /* overloads */ OverloadIndex(172),
},
{
/* [29] */
@@ -5647,13 +5684,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(171),
+ /* overloads */ OverloadIndex(173),
},
{
/* [32] */
/* fn vector_times_scalar[T : f32_f16, N : num](vec<N, T>, T) -> vec<N, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(172),
+ /* overloads */ OverloadIndex(174),
},
{
/* [33] */
@@ -5756,15 +5793,22 @@
},
{
/* [47] */
- /* fn sdot(u32, u32, u32) -> i32 */
- /* num overloads */ 1,
- /* overloads */ OverloadIndex(173),
+ /* fn faceForward[T : f32_f16](T, T, T) -> T */
+ /* fn faceForward[T : f32_f16, N : num](vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(162),
},
{
/* [48] */
+ /* fn sdot(u32, u32, u32) -> i32 */
+ /* num overloads */ 1,
+ /* overloads */ OverloadIndex(175),
+ },
+ {
+ /* [49] */
/* fn udot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(174),
+ /* overloads */ OverloadIndex(176),
},
};
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 ee7a65f..42f7461 100644
--- a/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
+++ b/src/tint/lang/spirv/reader/import_glsl_std450_test.cc
@@ -780,43 +780,6 @@
{"UnpackUnorm2x16", "unpack2x16unorm", 2},
{"UnpackHalf2x16", "unpack2x16float", 2}}));
-TEST_F(SpirvReaderTest, DISABLED_GlslStd450_FaceForward_Scalar) {
- // The %99 sum only has one use. Ensure it is evaluated only once by
- // making a let-declaration for it, since it is the normal operand to
- // the builtin function, and code generation uses it twice.
- EXPECT_IR(Preamble() + R"(
- %99 = OpFAdd %float %float_50 %float_50 ; normal operand has only one use
- %1 = OpExtInst %float %glsl FaceForward %99 %float_60 %float_70
- OpReturn
- OpFunctionEnd
- )",
- R"(
-%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
- $B1: {
- let x_1 = select(-(x_99), x_99, ((f2 * f3) < 0.0f));
- }
-}
-)");
-}
-
-TEST_F(SpirvReaderTest, DISABLED_GlslStd450_FaceForward_Vector) {
- EXPECT_IR(Preamble() + R"(
- %1 = OpExtInst %v2float %glsl FaceForward %v2float_50_60 %v2float_60_50 %v2float_70_70
- %2 = OpCopyObject %v2float %1
- OpReturn
- OpFunctionEnd
- )",
- R"(
-%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
- $B1: {
- %2:vec2<f32> = faceForward vec2<f32>(50.0f, 60.0f), vec2<f32>(60.0f, 50.0f), vec2<f32>(70.0f)
- %2:vec2<f32> = let %2
- ret
- }
-}
-)");
-}
-
TEST_F(SpirvReaderTest, DISABLED_GlslStd450_Reflect_Scalar) {
EXPECT_IR(Preamble() + R"(
%98 = OpFAdd %float %float_50 %float_50 ; has only one use
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index 639bdb3..f55198e 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -103,6 +103,9 @@
case spirv::BuiltinFn::kRefract:
Refract(builtin);
break;
+ case spirv::BuiltinFn::kFaceForward:
+ FaceForward(builtin);
+ break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
}
@@ -259,6 +262,25 @@
call->Destroy();
}
+ void FaceForward(spirv::ir::BuiltinCall* call) {
+ auto args = call->Args();
+ auto* N = args[0];
+ auto* I = args[1];
+ auto* Nref = args[2];
+
+ b.InsertBefore(call, [&] {
+ if (I->Type()->IsFloatScalar()) {
+ auto* neg = b.Negation(N->Type(), N);
+ auto* sel = b.Multiply(I->Type(), I, Nref)->Result(0);
+ sel = b.LessThan(ty.bool_(), sel, b.Zero(sel->Type()))->Result(0);
+ b.CallWithResult(call->DetachResult(), core::BuiltinFn::kSelect, neg, N, sel);
+ } else {
+ b.CallWithResult(call->DetachResult(), core::BuiltinFn::kFaceForward, N, I, Nref);
+ }
+ });
+ call->Destroy();
+ }
+
void Inverse(spirv::ir::BuiltinCall* call) {
auto* arg = call->Args()[0];
auto* mat_ty = arg->Type()->As<core::type::Matrix>();
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index 0fd0111..07cf354 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -2342,5 +2342,73 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvParser_BuiltinsTest, FaceForward_Scalar) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.f32(), spirv::BuiltinFn::kFaceForward, 50_f, 60_f, 70_f);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:f32 = spirv.faceForward 50.0f, 60.0f, 70.0f
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:f32 = negation 50.0f
+ %3:f32 = mul 60.0f, 70.0f
+ %4:bool = lt %3, 0.0f
+ %5:f32 = select %2, 50.0f, %4
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvParser_BuiltinsTest, FaceForward_Vector) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec2<f32>(), spirv::BuiltinFn::kFaceForward,
+ b.Splat(ty.vec2<f32>(), 10_f), b.Splat(ty.vec2<f32>(), 20_f),
+ b.Splat(ty.vec2<f32>(), 30_f));
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<f32> = spirv.faceForward vec2<f32>(10.0f), vec2<f32>(20.0f), vec2<f32>(30.0f)
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<f32> = faceForward vec2<f32>(10.0f), vec2<f32>(20.0f), vec2<f32>(30.0f)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/parser/import_glsl_std450_test.cc b/src/tint/lang/spirv/reader/parser/import_glsl_std450_test.cc
index 9e84be1..650464e 100644
--- a/src/tint/lang/spirv/reader/parser/import_glsl_std450_test.cc
+++ b/src/tint/lang/spirv/reader/parser/import_glsl_std450_test.cc
@@ -72,6 +72,7 @@
%v2float_50_60 = OpConstantComposite %v2float %float_50 %float_60
%v2float_60_50 = OpConstantComposite %v2float %float_60 %float_50
+ %v2float_70_60 = OpConstantComposite %v2float %float_70 %float_60
%v3float_50_60_70 = OpConstantComposite %v3float %float_50 %float_60 %float_70
%v4float_50_50_50_50 = OpConstantComposite %v4float %float_50 %float_50 %float_50 %float_50
@@ -583,5 +584,41 @@
)");
}
+TEST_F(SpirvParserTest, FaceForward_Scalar) {
+ EXPECT_IR(Preamble() + R"(
+ %99 = OpFAdd %float %float_50 %float_50 ; normal operand has only one use
+ %1 = OpExtInst %float %glsl FaceForward %99 %float_60 %float_70
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:f32 = add 50.0f, 50.0f
+ %3:f32 = spirv.faceForward %2, 60.0f, 70.0f
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, FaceForward_Vector) {
+ EXPECT_IR(Preamble() + R"(
+ %1 = OpExtInst %v2float %glsl FaceForward %v2float_50_60 %v2float_60_50 %v2float_70_60
+ %2 = OpCopyObject %v2float %1
+ OpReturn
+ OpFunctionEnd
+ )",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<f32> = spirv.faceForward vec2<f32>(50.0f, 60.0f), vec2<f32>(60.0f, 50.0f), vec2<f32>(70.0f, 60.0f)
+ %3:vec2<f32> = let %2
+ ret
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index b6b1ad3..bd1ad91 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -781,6 +781,8 @@
return spirv::BuiltinFn::kFindUMsb;
case GLSLstd450Refract:
return spirv::BuiltinFn::kRefract;
+ case GLSLstd450FaceForward:
+ return spirv::BuiltinFn::kFaceForward;
default:
break;
}
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index d2f16ac..1282a6d 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -355,6 +355,9 @@
implicit(T: f32_f16) fn refract(T, T, T) -> T
implicit(T: f32_f16, N: num) fn refract(vec<N, T>, vec<N, T>, T) -> vec<N, T>
+implicit(T: f32_f16) fn faceForward(T, T, T) -> T
+implicit(T: f32_f16, N: num) fn faceForward(vec<N, T>, vec<N, T>, vec<N, T>) -> vec<N, T>
+
////////////////////////////////////////////////////////////////////////////////
// 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 6df0a78..b10c68f 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1429,6 +1429,9 @@
case spirv::BuiltinFn::kRefract:
ext_inst(GLSLstd450Refract);
break;
+ case spirv::BuiltinFn::kFaceForward:
+ ext_inst(GLSLstd450FaceForward);
+ break;
case spirv::BuiltinFn::kUdot:
module_.PushExtension("SPV_KHR_integer_dot_product");
module_.PushCapability(SpvCapabilityDotProductKHR);
