[spirv-reader][ir] Add Signed comparison operators.
Add conversions for the `OpSGreaterThan`, `OpSGreaterThanEqual`,
`OpSLessThan` and `OpSLessThanEqual` comparison operators.
Bug: 391486024, 391486212, 391486712, 391486323
Change-Id: Id51db4896448dacb01710264dcdf6a048d3b2cc5
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/227214
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/spirv/builtin_fn.cc b/src/tint/lang/spirv/builtin_fn.cc
index 44edd7f..12bb13e 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -164,6 +164,14 @@
return "s_div";
case BuiltinFn::kSMod:
return "s_mod";
+ case BuiltinFn::kSGreaterThan:
+ return "s_greater_than";
+ case BuiltinFn::kSGreaterThanEqual:
+ return "s_greater_than_equal";
+ case BuiltinFn::kSLessThan:
+ return "s_less_than";
+ case BuiltinFn::kSLessThanEqual:
+ return "s_less_than_equal";
case BuiltinFn::kConvertFToS:
return "convertFToS";
case BuiltinFn::kConvertSToF:
@@ -278,6 +286,10 @@
case BuiltinFn::kBitwiseXor:
case BuiltinFn::kEqual:
case BuiltinFn::kNotEqual:
+ case BuiltinFn::kSGreaterThan:
+ case BuiltinFn::kSGreaterThanEqual:
+ case BuiltinFn::kSLessThan:
+ case BuiltinFn::kSLessThanEqual:
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 c99ba89..4e57ac1 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -111,6 +111,10 @@
case BuiltinFn::kBitwiseXor:
case BuiltinFn::kEqual:
case BuiltinFn::kNotEqual:
+ case BuiltinFn::kSGreaterThan:
+ case BuiltinFn::kSGreaterThanEqual:
+ case BuiltinFn::kSLessThan:
+ case BuiltinFn::kSLessThanEqual:
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 02778de..f643687 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -109,6 +109,10 @@
kMul,
kSDiv,
kSMod,
+ kSGreaterThan,
+ kSGreaterThanEqual,
+ kSLessThan,
+ kSLessThanEqual,
kConvertFToS,
kConvertSToF,
kConvertUToF,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 0b5486f..fe1a81a 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -5949,23 +5949,23 @@
{
/* [178] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 1,
+ /* num_parameters */ 2,
+ /* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(79),
- /* parameters */ ParameterIndex(1),
- /* return_matcher_indices */ MatcherIndicesIndex(4),
+ /* templates */ TemplateIndex(22),
+ /* parameters */ ParameterIndex(6),
+ /* return_matcher_indices */ MatcherIndicesIndex(17),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
/* [179] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 1,
- /* num_explicit_templates */ 1,
+ /* num_parameters */ 2,
+ /* num_explicit_templates */ 0,
/* num_templates */ 3,
- /* templates */ TemplateIndex(79),
- /* parameters */ ParameterIndex(362),
- /* return_matcher_indices */ MatcherIndicesIndex(149),
+ /* templates */ TemplateIndex(52),
+ /* parameters */ ParameterIndex(361),
+ /* return_matcher_indices */ MatcherIndicesIndex(179),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -5974,7 +5974,7 @@
/* num_parameters */ 1,
/* num_explicit_templates */ 1,
/* num_templates */ 2,
- /* templates */ TemplateIndex(82),
+ /* templates */ TemplateIndex(79),
/* parameters */ ParameterIndex(1),
/* return_matcher_indices */ MatcherIndicesIndex(4),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5985,7 +5985,7 @@
/* num_parameters */ 1,
/* num_explicit_templates */ 1,
/* num_templates */ 3,
- /* templates */ TemplateIndex(82),
+ /* templates */ TemplateIndex(79),
/* parameters */ ParameterIndex(362),
/* return_matcher_indices */ MatcherIndicesIndex(149),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5993,23 +5993,23 @@
{
/* [182] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 2,
- /* num_explicit_templates */ 0,
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 1,
/* num_templates */ 2,
- /* templates */ TemplateIndex(22),
- /* parameters */ ParameterIndex(6),
- /* return_matcher_indices */ MatcherIndicesIndex(17),
+ /* templates */ TemplateIndex(82),
+ /* parameters */ ParameterIndex(1),
+ /* return_matcher_indices */ MatcherIndicesIndex(4),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
/* [183] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
- /* num_parameters */ 2,
- /* num_explicit_templates */ 0,
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 1,
/* num_templates */ 3,
- /* templates */ TemplateIndex(52),
- /* parameters */ ParameterIndex(361),
- /* return_matcher_indices */ MatcherIndicesIndex(179),
+ /* templates */ TemplateIndex(82),
+ /* parameters */ ParameterIndex(362),
+ /* return_matcher_indices */ MatcherIndicesIndex(149),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
@@ -6727,86 +6727,114 @@
},
{
/* [61] */
- /* fn convertFToS<R : iu32>[T : f32_f16](T) -> R */
- /* fn convertFToS<R : iu32>[T : f32_f16, N : num](vec<N, T>) -> vec<N, R> */
+ /* fn s_greater_than[A : iu32, B : iu32](A, B) -> bool */
+ /* fn s_greater_than[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, bool> */
/* num overloads */ 2,
/* overloads */ OverloadIndex(178),
},
{
/* [62] */
- /* fn convertSToF<R : f32_f16>[T : iu32](T) -> R */
- /* fn convertSToF<R : f32_f16>[T : iu32, N : num](vec<N, T>) -> vec<N, R> */
+ /* fn s_greater_than_equal[A : iu32, B : iu32](A, B) -> bool */
+ /* fn s_greater_than_equal[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, bool> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(180),
+ /* overloads */ OverloadIndex(178),
},
{
/* [63] */
- /* fn convertUToF<R : f32_f16>[T : iu32](T) -> R */
- /* fn convertUToF<R : f32_f16>[T : iu32, N : num](vec<N, T>) -> vec<N, R> */
+ /* fn s_less_than[A : iu32, B : iu32](A, B) -> bool */
+ /* fn s_less_than[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, bool> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(178),
+ },
+ {
+ /* [64] */
+ /* fn s_less_than_equal[A : iu32, B : iu32](A, B) -> bool */
+ /* fn s_less_than_equal[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, bool> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(178),
+ },
+ {
+ /* [65] */
+ /* fn convertFToS<R : iu32>[T : f32_f16](T) -> R */
+ /* fn convertFToS<R : iu32>[T : f32_f16, N : num](vec<N, T>) -> vec<N, R> */
/* num overloads */ 2,
/* overloads */ OverloadIndex(180),
},
{
- /* [64] */
+ /* [66] */
+ /* fn convertSToF<R : f32_f16>[T : iu32](T) -> R */
+ /* fn convertSToF<R : f32_f16>[T : iu32, N : num](vec<N, T>) -> vec<N, R> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(182),
+ },
+ {
+ /* [67] */
+ /* fn convertUToF<R : f32_f16>[T : iu32](T) -> R */
+ /* fn convertUToF<R : f32_f16>[T : iu32, N : num](vec<N, T>) -> vec<N, R> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(182),
+ },
+ {
+ /* [68] */
/* fn bitwiseAnd<R : iu32>[A : iu32, B : iu32](A, B) -> R */
/* fn bitwiseAnd<R : iu32>[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, R> */
/* num overloads */ 2,
/* overloads */ OverloadIndex(176),
},
{
- /* [65] */
+ /* [69] */
/* fn bitwiseOr<R : iu32>[A : iu32, B : iu32](A, B) -> R */
/* fn bitwiseOr<R : iu32>[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, R> */
/* num overloads */ 2,
/* overloads */ OverloadIndex(176),
},
{
- /* [66] */
+ /* [70] */
/* fn bitwiseXor<R : iu32>[A : iu32, B : iu32](A, B) -> R */
/* fn bitwiseXor<R : iu32>[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, R> */
/* num overloads */ 2,
/* overloads */ OverloadIndex(176),
},
{
- /* [67] */
+ /* [71] */
/* fn equal[A : iu32, B : iu32](A, B) -> bool */
/* fn equal[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, bool> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(182),
+ /* overloads */ OverloadIndex(178),
},
{
- /* [68] */
+ /* [72] */
/* fn not_equal[A : iu32, B : iu32](A, B) -> bool */
/* fn not_equal[A : iu32, B : iu32, N : num](vec<N, A>, vec<N, B>) -> vec<N, bool> */
/* num overloads */ 2,
- /* overloads */ OverloadIndex(182),
+ /* overloads */ OverloadIndex(178),
},
{
- /* [69] */
+ /* [73] */
/* fn sdot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(195),
},
{
- /* [70] */
+ /* [74] */
/* fn udot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(196),
},
{
- /* [71] */
+ /* [75] */
/* fn cooperative_matrix_load<T : subgroup_matrix<K, S, C, R>>[K : subgroup_matrix_kind, S : fiu32_f16, C : num, R : num](ptr<workgroup_or_storage, S, readable>, u32, u32, u32) -> T */
/* num overloads */ 1,
/* overloads */ OverloadIndex(197),
},
{
- /* [72] */
+ /* [76] */
/* fn cooperative_matrix_store[K : subgroup_matrix_kind, S : fiu32_f16, C : num, R : num](ptr<workgroup_or_storage, S, writable>, subgroup_matrix<K, S, C, R>, u32, u32, u32) */
/* num overloads */ 1,
/* overloads */ OverloadIndex(198),
},
{
- /* [73] */
+ /* [77] */
/* fn cooperative_matrix_mul_add[T : subgroup_matrix_elements, TR : subgroup_matrix_elements, C : num, R : num, K : num](subgroup_matrix<subgroup_matrix_kind_left, T, K, R>, subgroup_matrix<subgroup_matrix_kind_right, T, C, K>, subgroup_matrix<subgroup_matrix_kind_result, TR, C, R>, u32) -> subgroup_matrix<subgroup_matrix_kind_result, TR, C, R> */
/* num overloads */ 1,
/* overloads */ OverloadIndex(199),
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index 170385c..e4fee04 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -170,6 +170,18 @@
case spirv::BuiltinFn::kNotEqual:
NotEqual(builtin);
break;
+ case spirv::BuiltinFn::kSGreaterThan:
+ SGreaterThan(builtin);
+ break;
+ case spirv::BuiltinFn::kSGreaterThanEqual:
+ SGreaterThanEqual(builtin);
+ break;
+ case spirv::BuiltinFn::kSLessThan:
+ SLessThan(builtin);
+ break;
+ case spirv::BuiltinFn::kSLessThanEqual:
+ SLessThanEqual(builtin);
+ break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
}
@@ -313,6 +325,37 @@
EmitBinaryMatchedArgs(call, core::BinaryOp::kNotEqual);
}
+ void EmitBinaryWithSignedArgs(spirv::ir::BuiltinCall* call, core::BinaryOp op) {
+ const auto& args = call->Args();
+ auto* lhs = args[0];
+ auto* rhs = args[1];
+
+ auto* arg_ty = ty.MatchWidth(ty.i32(), call->Result(0)->Type());
+ b.InsertBefore(call, [&] {
+ if (lhs->Type() != arg_ty) {
+ lhs = b.Bitcast(arg_ty, lhs)->Result(0);
+ }
+ if (rhs->Type() != arg_ty) {
+ rhs = b.Bitcast(arg_ty, rhs)->Result(0);
+ }
+
+ b.BinaryWithResult(call->DetachResult(), op, lhs, rhs)->Result(0);
+ });
+ call->Destroy();
+ }
+ void SGreaterThan(spirv::ir::BuiltinCall* call) {
+ EmitBinaryWithSignedArgs(call, core::BinaryOp::kGreaterThan);
+ }
+ void SGreaterThanEqual(spirv::ir::BuiltinCall* call) {
+ EmitBinaryWithSignedArgs(call, core::BinaryOp::kGreaterThanEqual);
+ }
+ void SLessThan(spirv::ir::BuiltinCall* call) {
+ EmitBinaryWithSignedArgs(call, core::BinaryOp::kLessThan);
+ }
+ void SLessThanEqual(spirv::ir::BuiltinCall* call) {
+ EmitBinaryWithSignedArgs(call, core::BinaryOp::kLessThanEqual);
+ }
+
// The SPIR-V Signed methods all interpret their arguments as signed (regardless of the type of
// the argument). In order to satisfy this, we must bitcast any unsigned argument to a signed
// type before calling the WGSL equivalent method.
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index 139e5e4..8cf5cff 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -6305,5 +6305,291 @@
::testing::Values(SpirvReaderParams{spirv::BuiltinFn::kEqual, "equal", "eq"},
SpirvReaderParams{spirv::BuiltinFn::kNotEqual, "not_equal", "neq"}));
+using SpirvReader_SignedIntegerTest =
+ core::ir::transform::TransformTestWithParam<SpirvReaderParams>;
+TEST_P(SpirvReader_SignedIntegerTest, Scalar_SignedSigned) {
+ auto param = GetParam();
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.bool_(), param.fn, 1_i, 2_i);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = spirv.)" +
+ param.spv_name + R"( 1i, 2i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = )" + param.wgsl_name +
+ R"( 1i, 2i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_P(SpirvReader_SignedIntegerTest, Scalar_SignedUnsigned) {
+ auto param = GetParam();
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.bool_(), param.fn, 1_i, 8_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = spirv.)" +
+ param.spv_name + R"( 1i, 8u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = bitcast 8u
+ %3:bool = )" + param.wgsl_name +
+ R"( 1i, %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_P(SpirvReader_SignedIntegerTest, Scalar_UnsignedSigned) {
+ auto param = GetParam();
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.bool_(), param.fn, 8_u, 1_i);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = spirv.)" +
+ param.spv_name + R"( 8u, 1i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = bitcast 8u
+ %3:bool = )" + param.wgsl_name +
+ R"( %2, 1i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_P(SpirvReader_SignedIntegerTest, Scalar_UnsignedUnsigned) {
+ auto param = GetParam();
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.bool_(), param.fn, 8_u, 9_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = spirv.)" +
+ param.spv_name + R"( 8u, 9u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = bitcast 8u
+ %3:i32 = bitcast 9u
+ %4:bool = )" + param.wgsl_name +
+ R"( %2, %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_P(SpirvReader_SignedIntegerTest, Vector_SignedSigned) {
+ auto param = GetParam();
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec2<bool>(), param.fn, b.Splat<vec2<i32>>(1_i),
+ b.Splat<vec2<i32>>(2_i));
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<bool> = spirv.)" +
+ param.spv_name + R"( vec2<i32>(1i), vec2<i32>(2i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<bool> = )" +
+ param.wgsl_name + R"( vec2<i32>(1i), vec2<i32>(2i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_P(SpirvReader_SignedIntegerTest, Vector_SignedUnsigned) {
+ auto param = GetParam();
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec2<bool>(), param.fn, b.Splat<vec2<i32>>(1_i),
+ b.Splat<vec2<u32>>(8_u));
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<bool> = spirv.)" +
+ param.spv_name + R"( vec2<i32>(1i), vec2<u32>(8u)
+ 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>(8u)
+ %3:vec2<bool> = )" +
+ param.wgsl_name + R"( vec2<i32>(1i), %2
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_P(SpirvReader_SignedIntegerTest, Vector_UnsignedSigned) {
+ auto param = GetParam();
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec2<bool>(), param.fn, b.Splat<vec2<u32>>(8_u),
+ b.Splat<vec2<i32>>(1_i));
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<bool> = spirv.)" +
+ param.spv_name + R"( vec2<u32>(8u), vec2<i32>(1i)
+ 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>(8u)
+ %3:vec2<bool> = )" +
+ param.wgsl_name + R"( %2, vec2<i32>(1i)
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+
+TEST_P(SpirvReader_SignedIntegerTest, Vector_UnsignedUnsigned) {
+ auto param = GetParam();
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.Call<spirv::ir::BuiltinCall>(ty.vec2<bool>(), param.fn, b.Splat<vec2<u32>>(8_u),
+ b.Splat<vec2<u32>>(9_u));
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<bool> = spirv.)" +
+ param.spv_name + R"( vec2<u32>(8u), vec2<u32>(9u)
+ 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>(8u)
+ %3:vec2<i32> = bitcast vec2<u32>(9u)
+ %4:vec2<bool> = )" +
+ param.wgsl_name + R"( %2, %3
+ ret
+ }
+}
+)";
+ EXPECT_EQ(expect, str());
+}
+INSTANTIATE_TEST_SUITE_P(
+ SpirvReader,
+ SpirvReader_SignedIntegerTest,
+ ::testing::Values(
+ SpirvReaderParams{spirv::BuiltinFn::kSGreaterThan, "s_greater_than", "gt"},
+ SpirvReaderParams{spirv::BuiltinFn::kSGreaterThanEqual, "s_greater_than_equal", "gte"},
+ SpirvReaderParams{spirv::BuiltinFn::kSLessThan, "s_less_than", "lt"},
+ SpirvReaderParams{spirv::BuiltinFn::kSLessThanEqual, "s_less_than_equal", "lte"}));
+
} // namespace
} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/parser/logical_test.cc b/src/tint/lang/spirv/reader/parser/logical_test.cc
index f00c282..a0becf7 100644
--- a/src/tint/lang/spirv/reader/parser/logical_test.cc
+++ b/src/tint/lang/spirv/reader/parser/logical_test.cc
@@ -232,7 +232,7 @@
%ep_type = OpTypeFunction %void
%main = OpFunction %void None %ep_type
%main_start = OpLabel
- %1 = OpI)" +
+ %1 = Op)" +
params.spv_name + R"( %bool %one %two
OpReturn
OpFunctionEnd
@@ -274,7 +274,7 @@
%ep_type = OpTypeFunction %void
%main = OpFunction %void None %ep_type
%main_start = OpLabel
- %1 = OpI)" +
+ %1 = Op)" +
params.spv_name + R"( %bool %one %eight
OpReturn
OpFunctionEnd
@@ -316,7 +316,7 @@
%ep_type = OpTypeFunction %void
%main = OpFunction %void None %ep_type
%main_start = OpLabel
- %1 = OpI)" +
+ %1 = Op)" +
params.spv_name + R"( %bool %eight %one
OpReturn
OpFunctionEnd
@@ -358,7 +358,7 @@
%ep_type = OpTypeFunction %void
%main = OpFunction %void None %ep_type
%main_start = OpLabel
- %1 = OpI)" +
+ %1 = Op)" +
params.spv_name + R"( %bool %eight %nine
OpReturn
OpFunctionEnd
@@ -400,7 +400,7 @@
%ep_type = OpTypeFunction %void
%main = OpFunction %void None %ep_type
%main_start = OpLabel
- %1 = OpI)" +
+ %1 = Op)" +
params.spv_name + R"( %v2bool %v2one %v2two
OpReturn
OpFunctionEnd
@@ -442,7 +442,7 @@
%ep_type = OpTypeFunction %void
%main = OpFunction %void None %ep_type
%main_start = OpLabel
- %1 = OpI)" +
+ %1 = Op)" +
params.spv_name + R"( %v2bool %v2one %v2eight
OpReturn
OpFunctionEnd
@@ -484,7 +484,7 @@
%ep_type = OpTypeFunction %void
%main = OpFunction %void None %ep_type
%main_start = OpLabel
- %1 = OpI)" +
+ %1 = Op)" +
params.spv_name + R"( %v2bool %v2eight %v2one
OpReturn
OpFunctionEnd
@@ -526,7 +526,7 @@
%ep_type = OpTypeFunction %void
%main = OpFunction %void None %ep_type
%main_start = OpLabel
- %1 = OpI)" +
+ %1 = Op)" +
params.spv_name + R"( %v2bool %v2eight %v2nine
OpReturn
OpFunctionEnd
@@ -544,8 +544,13 @@
INSTANTIATE_TEST_SUITE_P(SpirvParserTest,
SpirvParser_IntegerTest,
- testing::Values(SpirvLogicalParam{"Equal", "equal"},
- SpirvLogicalParam{"NotEqual", "not_equal"}));
+ testing::Values(SpirvLogicalParam{"IEqual", "equal"},
+ SpirvLogicalParam{"INotEqual", "not_equal"},
+ SpirvLogicalParam{"SGreaterThan", "s_greater_than"},
+ SpirvLogicalParam{"SGreaterThanEqual",
+ "s_greater_than_equal"},
+ SpirvLogicalParam{"SLessThan", "s_less_than"},
+ SpirvLogicalParam{"SLessThanEqual", "s_less_than_equal"}));
using SpirvParser_LogicalTest = SpirvParserTestWithParam<SpirvLogicalParam>;
TEST_P(SpirvParser_LogicalTest, Scalar) {
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index a9e2543..ed86efc 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -764,6 +764,18 @@
case spv::Op::OpINotEqual:
EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kNotEqual);
break;
+ case spv::Op::OpSGreaterThan:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kSGreaterThan);
+ break;
+ case spv::Op::OpSGreaterThanEqual:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kSGreaterThanEqual);
+ break;
+ case spv::Op::OpSLessThan:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kSLessThan);
+ break;
+ case spv::Op::OpSLessThanEqual:
+ EmitSpirvBuiltinCall(inst, spirv::BuiltinFn::kSLessThanEqual);
+ break;
case spv::Op::OpISub:
EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kSub);
break;
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index 2833eba..82ad82f 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -414,6 +414,15 @@
implicit(A: iu32, B: iu32, N: num) fn s_div<R: iu32>(vec<N, A>, vec<N, B>) -> vec<N, R>
implicit(A: iu32, B: iu32, N: num) fn s_mod<R: iu32>(vec<N, A>, vec<N, B>) -> vec<N, R>
+implicit(A: iu32, B: iu32) fn s_greater_than(A, B) -> bool
+implicit(A: iu32, B: iu32, N: num) fn s_greater_than(vec<N, A>, vec<N, B>) -> vec<N, bool>
+implicit(A: iu32, B: iu32) fn s_greater_than_equal(A, B) -> bool
+implicit(A: iu32, B: iu32, N: num) fn s_greater_than_equal(vec<N, A>, vec<N, B>) -> vec<N, bool>
+implicit(A: iu32, B: iu32) fn s_less_than(A, B) -> bool
+implicit(A: iu32, B: iu32, N: num) fn s_less_than(vec<N, A>, vec<N, B>) -> vec<N, bool>
+implicit(A: iu32, B: iu32) fn s_less_than_equal(A, B) -> bool
+implicit(A: iu32, B: iu32, N: num) fn s_less_than_equal(vec<N, A>, vec<N, B>) -> vec<N, bool>
+
implicit(T: f32_f16) fn convertFToS<R: iu32>(T) -> R
implicit(T: f32_f16, N: num) fn convertFToS<R: iu32>(vec<N, T>) -> vec<N, R>
implicit(T: iu32) fn convertSToF<R: f32_f16>(T) -> R
diff --git a/src/tint/lang/spirv/writer/printer/printer.cc b/src/tint/lang/spirv/writer/printer/printer.cc
index c76d552..c8d3f9e 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1548,6 +1548,18 @@
case BuiltinFn::kNotEqual:
op = spv::Op::OpINotEqual;
break;
+ case BuiltinFn::kSGreaterThan:
+ op = spv::Op::OpSGreaterThan;
+ break;
+ case BuiltinFn::kSGreaterThanEqual:
+ op = spv::Op::OpSGreaterThanEqual;
+ break;
+ case BuiltinFn::kSLessThan:
+ op = spv::Op::OpSLessThan;
+ break;
+ case BuiltinFn::kSLessThanEqual:
+ op = spv::Op::OpSLessThanEqual;
+ break;
case spirv::BuiltinFn::kNone:
TINT_ICE() << "undefined spirv ir function";
}
