[spirv-reader][ir] Add support for OpSNegate
Handle the `OpSNegate` SPIR-V instruction conversion to IR.
Bug: 391486172
Change-Id: Ia58c67b487797529b5b6a5390ec4fb351978f3a4
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/228114
Reviewed-by: James Price <jrprice@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 6543eca..4375efe 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -208,6 +208,8 @@
return "shift_right_arithmetic";
case BuiltinFn::kNot:
return "not";
+ case BuiltinFn::kSNegate:
+ return "s_negate";
case BuiltinFn::kSDot:
return "s_dot";
case BuiltinFn::kUDot:
@@ -320,6 +322,7 @@
case BuiltinFn::kShiftRightLogical:
case BuiltinFn::kShiftRightArithmetic:
case BuiltinFn::kNot:
+ case BuiltinFn::kSNegate:
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 b82b996..71f9828 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -125,6 +125,7 @@
case BuiltinFn::kShiftRightLogical:
case BuiltinFn::kShiftRightArithmetic:
case BuiltinFn::kNot:
+ case BuiltinFn::kSNegate:
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 0c6d59b..9b78b05 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -131,6 +131,7 @@
kShiftRightLogical,
kShiftRightArithmetic,
kNot,
+ kSNegate,
kSDot,
kUDot,
kCooperativeMatrixLoad,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 1eea028..09c40ea 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -7034,30 +7034,37 @@
},
{
/* [83] */
+ /* fn s_negate<R : iu32>[A : iu32](A) -> R */
+ /* fn s_negate<R : iu32>[A : iu32, N : num](vec<N, A>) -> vec<N, R> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(192),
+ },
+ {
+ /* [84] */
/* fn s_dot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(202),
},
{
- /* [84] */
+ /* [85] */
/* fn u_dot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(203),
},
{
- /* [85] */
+ /* [86] */
/* 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(204),
},
{
- /* [86] */
+ /* [87] */
/* 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(205),
},
{
- /* [87] */
+ /* [88] */
/* 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(206),
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index f267065..1dc533f 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -206,12 +206,34 @@
case spirv::BuiltinFn::kNot:
Not(builtin);
break;
+ case spirv::BuiltinFn::kSNegate:
+ SNegate(builtin);
+ break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
}
}
}
+ void SNegate(spirv::ir::BuiltinCall* call) {
+ auto* val = call->Args()[0];
+
+ auto* res_ty = call->Result(0)->Type();
+ auto* neg_ty = ty.MatchWidth(ty.i32(), val->Type());
+ b.InsertBefore(call, [&] {
+ if (val->Type() != neg_ty) {
+ val = b.Bitcast(neg_ty, val)->Result(0);
+ }
+ val = b.Negation(neg_ty, val)->Result(0);
+
+ if (neg_ty != res_ty) {
+ val = b.Bitcast(res_ty, val)->Result(0);
+ }
+ call->Result(0)->ReplaceAllUsesWith(val);
+ });
+ call->Destroy();
+ }
+
void Not(spirv::ir::BuiltinCall* call) {
auto* val = call->Args()[0];
auto* result_ty = call->Result(0)->Type();
diff --git a/src/tint/lang/spirv/reader/lower/builtins_test.cc b/src/tint/lang/spirv/reader/lower/builtins_test.cc
index 91d6f1f..9457d54 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -8729,5 +8729,269 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvReader_BuiltinsTest, SNegate_Scalar_Signed_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kSNegate,
+ Vector{ty.i32()}, 1_i);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.s_negate<i32> 1i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = negation 1i
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, SNegate_Scalar_Signed_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kSNegate,
+ Vector{ty.u32()}, 1_i);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.s_negate<u32> 1i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = negation 1i
+ %3:u32 = bitcast %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, SNegate_Scalar_Unsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kSNegate,
+ Vector{ty.i32()}, 8_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.s_negate<i32> 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:i32 = negation %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, SNegate_Scalar_Unsigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kSNegate,
+ Vector{ty.u32()}, 8_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.s_negate<u32> 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:i32 = negation %2
+ %4:u32 = bitcast %3
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, SNegate_Vector_Signed_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kSNegate,
+ Vector{ty.i32()}, b.Splat<vec2<i32>>(1_i));
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.s_negate<i32> 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> = negation vec2<i32>(1i)
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, SNegate_Vector_Signed_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kSNegate,
+ Vector{ty.u32()}, b.Splat<vec2<i32>>(1_i));
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.s_negate<u32> 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> = negation vec2<i32>(1i)
+ %3:vec2<u32> = bitcast %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, SNegate_Vector_Unsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kSNegate,
+ Vector{ty.i32()}, b.Splat<vec2<u32>>(8_u));
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.s_negate<i32> 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<i32> = negation %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, SNegate_Vector_Unsigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kSNegate,
+ Vector{ty.u32()}, b.Splat<vec2<u32>>(8_u));
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.s_negate<u32> 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<i32> = negation %2
+ %4:vec2<u32> = bitcast %3
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::spirv::reader::lower
diff --git a/src/tint/lang/spirv/reader/parser/bit_test.cc b/src/tint/lang/spirv/reader/parser/bit_test.cc
index df30fec..eb6004a 100644
--- a/src/tint/lang/spirv/reader/parser/bit_test.cc
+++ b/src/tint/lang/spirv/reader/parser/bit_test.cc
@@ -1654,5 +1654,277 @@
)");
}
+TEST_F(SpirvParserTest, SNegate_Scalar_Signed_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2int = OpTypeVector %int 2
+ %v2uint = OpTypeVector %uint 2
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpSNegate %int %one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.s_negate<i32> 1i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, SNegate_Scalar_Signed_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2int = OpTypeVector %int 2
+ %v2uint = OpTypeVector %uint 2
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpSNegate %uint %one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.s_negate<u32> 1i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, SNegate_Scalar_Unsigned_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2int = OpTypeVector %int 2
+ %v2uint = OpTypeVector %uint 2
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpSNegate %int %eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.s_negate<i32> 8u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, SNegate_Scalar_Unsigned_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2int = OpTypeVector %int 2
+ %v2uint = OpTypeVector %uint 2
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpSNegate %uint %eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.s_negate<u32> 8u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, SNegate_Vector_Signed_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2int = OpTypeVector %int 2
+ %v2uint = OpTypeVector %uint 2
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpSNegate %v2int %v2one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.s_negate<i32> vec2<i32>(1i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, SNegate_Vector_Signed_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2int = OpTypeVector %int 2
+ %v2uint = OpTypeVector %uint 2
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpSNegate %v2uint %v2one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.s_negate<u32> vec2<i32>(1i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, SNegate_Vector_Unsigned_Signed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2int = OpTypeVector %int 2
+ %v2uint = OpTypeVector %uint 2
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpSNegate %v2int %v2eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.s_negate<i32> vec2<u32>(8u)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, SNegate_Vector_Unsigned_Unsigned) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %int = OpTypeInt 32 1
+ %uint = OpTypeInt 32 0
+ %one = OpConstant %int 1
+ %two = OpConstant %int 2
+ %eight = OpConstant %uint 8
+ %nine = OpConstant %uint 9
+ %v2int = OpTypeVector %int 2
+ %v2uint = OpTypeVector %uint 2
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpSNegate %v2uint %v2eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.s_negate<u32> vec2<u32>(8u)
+ 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 e06fd0e..adcc929 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -979,6 +979,9 @@
case spv::Op::OpTranspose:
EmitBuiltinCall(inst, core::BuiltinFn::kTranspose);
break;
+ case spv::Op::OpSNegate:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kSNegate);
+ break;
default:
TINT_UNIMPLEMENTED()
<< "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index 3feb5fa..e11cce8 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -480,6 +480,9 @@
implicit(A: iu32) fn not<R: iu32>(A) -> R
implicit(A: iu32, N: num) fn not<R: iu32>(vec<N, A>) -> vec<N, R>
+implicit(A: iu32) fn s_negate<R: iu32>(A) -> R
+implicit(A: iu32, N: num) fn s_negate<R: iu32>(vec<N, A>) -> 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 eb55006..a611c3b 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1590,6 +1590,9 @@
case BuiltinFn::kNot:
op = spv::Op::OpNot;
break;
+ case BuiltinFn::kSNegate:
+ op = spv::Op::OpSNegate;
+ break;
case spirv::BuiltinFn::kNone:
TINT_ICE() << "undefined spirv ir function";
}
