[spirv-reader][ir] Convert OpNot.
This CL converts the SPIR-V `OpNot` instruction into an IR `complement`
instruction.
Bug: 391487408
Change-Id: I6c2b5745b8c03bd513c7a0237eaa38693611aa4a
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/227974
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 80c038d..6543eca 100644
--- a/src/tint/lang/spirv/builtin_fn.cc
+++ b/src/tint/lang/spirv/builtin_fn.cc
@@ -206,6 +206,8 @@
return "shift_right_logical";
case BuiltinFn::kShiftRightArithmetic:
return "shift_right_arithmetic";
+ case BuiltinFn::kNot:
+ return "not";
case BuiltinFn::kSDot:
return "s_dot";
case BuiltinFn::kUDot:
@@ -317,6 +319,7 @@
case BuiltinFn::kShiftLeftLogical:
case BuiltinFn::kShiftRightLogical:
case BuiltinFn::kShiftRightArithmetic:
+ case BuiltinFn::kNot:
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 4eeafca..b82b996 100644
--- a/src/tint/lang/spirv/builtin_fn.cc.tmpl
+++ b/src/tint/lang/spirv/builtin_fn.cc.tmpl
@@ -124,6 +124,7 @@
case BuiltinFn::kShiftLeftLogical:
case BuiltinFn::kShiftRightLogical:
case BuiltinFn::kShiftRightArithmetic:
+ case BuiltinFn::kNot:
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 d37582c..0c6d59b 100644
--- a/src/tint/lang/spirv/builtin_fn.h
+++ b/src/tint/lang/spirv/builtin_fn.h
@@ -130,6 +130,7 @@
kShiftLeftLogical,
kShiftRightLogical,
kShiftRightArithmetic,
+ kNot,
kSDot,
kUDot,
kCooperativeMatrixLoad,
diff --git a/src/tint/lang/spirv/intrinsic/data.cc b/src/tint/lang/spirv/intrinsic/data.cc
index 3dca8fb..1eea028 100644
--- a/src/tint/lang/spirv/intrinsic/data.cc
+++ b/src/tint/lang/spirv/intrinsic/data.cc
@@ -3983,48 +3983,66 @@
},
{
/* [85] */
- /* name */ "I",
- /* matcher_indices */ MatcherIndicesIndex(127),
+ /* name */ "R",
+ /* matcher_indices */ MatcherIndicesIndex(234),
/* kind */ TemplateInfo::Kind::kType,
},
{
/* [86] */
/* name */ "A",
- /* matcher_indices */ MatcherIndicesIndex(/* invalid */),
- /* kind */ TemplateInfo::Kind::kNumber,
+ /* matcher_indices */ MatcherIndicesIndex(234),
+ /* kind */ TemplateInfo::Kind::kType,
},
{
/* [87] */
- /* name */ "F",
+ /* name */ "N",
/* matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* kind */ TemplateInfo::Kind::kNumber,
},
{
/* [88] */
+ /* name */ "I",
+ /* matcher_indices */ MatcherIndicesIndex(127),
+ /* kind */ TemplateInfo::Kind::kType,
+ },
+ {
+ /* [89] */
/* name */ "A",
/* matcher_indices */ MatcherIndicesIndex(/* invalid */),
/* kind */ TemplateInfo::Kind::kNumber,
},
{
- /* [89] */
+ /* [90] */
+ /* name */ "F",
+ /* matcher_indices */ MatcherIndicesIndex(/* invalid */),
+ /* kind */ TemplateInfo::Kind::kNumber,
+ },
+ {
+ /* [91] */
+ /* name */ "A",
+ /* matcher_indices */ MatcherIndicesIndex(/* invalid */),
+ /* kind */ TemplateInfo::Kind::kNumber,
+ },
+ {
+ /* [92] */
/* name */ "T",
/* matcher_indices */ MatcherIndicesIndex(236),
/* kind */ TemplateInfo::Kind::kType,
},
{
- /* [90] */
+ /* [93] */
/* name */ "S",
/* matcher_indices */ MatcherIndicesIndex(238),
/* kind */ TemplateInfo::Kind::kType,
},
{
- /* [91] */
+ /* [94] */
/* name */ "T",
/* matcher_indices */ MatcherIndicesIndex(235),
/* kind */ TemplateInfo::Kind::kType,
},
{
- /* [92] */
+ /* [95] */
/* name */ "S",
/* matcher_indices */ MatcherIndicesIndex(240),
/* kind */ TemplateInfo::Kind::kNumber,
@@ -4679,7 +4697,7 @@
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(87),
+ /* templates */ TemplateIndex(90),
/* parameters */ ParameterIndex(293),
/* return_matcher_indices */ MatcherIndicesIndex(127),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -4690,7 +4708,7 @@
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(87),
+ /* templates */ TemplateIndex(90),
/* parameters */ ParameterIndex(296),
/* return_matcher_indices */ MatcherIndicesIndex(214),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -4701,7 +4719,7 @@
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(87),
+ /* templates */ TemplateIndex(90),
/* parameters */ ParameterIndex(299),
/* return_matcher_indices */ MatcherIndicesIndex(216),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -4712,7 +4730,7 @@
/* num_parameters */ 1,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(87),
+ /* templates */ TemplateIndex(90),
/* parameters */ ParameterIndex(302),
/* return_matcher_indices */ MatcherIndicesIndex(216),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5262,7 +5280,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(89),
+ /* templates */ TemplateIndex(92),
/* parameters */ ParameterIndex(142),
/* return_matcher_indices */ MatcherIndicesIndex(93),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5273,7 +5291,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(89),
+ /* templates */ TemplateIndex(92),
/* parameters */ ParameterIndex(333),
/* return_matcher_indices */ MatcherIndicesIndex(96),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5284,7 +5302,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(89),
+ /* templates */ TemplateIndex(92),
/* parameters */ ParameterIndex(335),
/* return_matcher_indices */ MatcherIndicesIndex(99),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5295,7 +5313,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(89),
+ /* templates */ TemplateIndex(92),
/* parameters */ ParameterIndex(337),
/* return_matcher_indices */ MatcherIndicesIndex(102),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5306,7 +5324,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(89),
+ /* templates */ TemplateIndex(92),
/* parameters */ ParameterIndex(339),
/* return_matcher_indices */ MatcherIndicesIndex(105),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5317,7 +5335,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(89),
+ /* templates */ TemplateIndex(92),
/* parameters */ ParameterIndex(341),
/* return_matcher_indices */ MatcherIndicesIndex(108),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5328,7 +5346,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 1,
- /* templates */ TemplateIndex(90),
+ /* templates */ TemplateIndex(93),
/* parameters */ ParameterIndex(343),
/* return_matcher_indices */ MatcherIndicesIndex(192),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5339,7 +5357,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 1,
- /* templates */ TemplateIndex(90),
+ /* templates */ TemplateIndex(93),
/* parameters */ ParameterIndex(345),
/* return_matcher_indices */ MatcherIndicesIndex(198),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5350,7 +5368,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 1,
- /* templates */ TemplateIndex(90),
+ /* templates */ TemplateIndex(93),
/* parameters */ ParameterIndex(347),
/* return_matcher_indices */ MatcherIndicesIndex(200),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5361,7 +5379,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 1,
- /* templates */ TemplateIndex(90),
+ /* templates */ TemplateIndex(93),
/* parameters */ ParameterIndex(349),
/* return_matcher_indices */ MatcherIndicesIndex(202),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -5977,7 +5995,7 @@
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(91),
+ /* templates */ TemplateIndex(94),
/* parameters */ ParameterIndex(361),
/* return_matcher_indices */ MatcherIndicesIndex(4),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
@@ -6150,16 +6168,38 @@
{
/* [192] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 1,
+ /* num_templates */ 2,
+ /* templates */ TemplateIndex(51),
+ /* parameters */ ParameterIndex(1),
+ /* return_matcher_indices */ MatcherIndicesIndex(4),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [193] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+ /* num_parameters */ 1,
+ /* num_explicit_templates */ 1,
+ /* num_templates */ 3,
+ /* templates */ TemplateIndex(85),
+ /* parameters */ ParameterIndex(370),
+ /* return_matcher_indices */ MatcherIndicesIndex(159),
+ /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+ },
+ {
+ /* [194] */
+ /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
/* num_templates */ 2,
- /* templates */ TemplateIndex(85),
+ /* templates */ TemplateIndex(88),
/* parameters */ ParameterIndex(331),
/* return_matcher_indices */ MatcherIndicesIndex(127),
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [193] */
+ /* [195] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -6170,7 +6210,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [194] */
+ /* [196] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -6181,7 +6221,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [195] */
+ /* [197] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -6192,7 +6232,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [196] */
+ /* [198] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -6203,7 +6243,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [197] */
+ /* [199] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -6214,7 +6254,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [198] */
+ /* [200] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -6225,7 +6265,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [199] */
+ /* [201] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 2,
/* num_explicit_templates */ 0,
@@ -6236,7 +6276,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [200] */
+ /* [202] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -6247,7 +6287,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [201] */
+ /* [203] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 3,
/* num_explicit_templates */ 0,
@@ -6258,7 +6298,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [202] */
+ /* [204] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 4,
/* num_explicit_templates */ 1,
@@ -6269,7 +6309,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [203] */
+ /* [205] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
/* num_parameters */ 5,
/* num_explicit_templates */ 0,
@@ -6280,7 +6320,7 @@
/* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
},
{
- /* [204] */
+ /* [206] */
/* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMustUse),
/* num_parameters */ 4,
/* num_explicit_templates */ 0,
@@ -6300,7 +6340,7 @@
/* [0] */
/* fn array_length[I : u32, A : access](ptr<storage, struct_with_runtime_array, A>, I) -> u32 */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(192),
+ /* overloads */ OverloadIndex(194),
},
{
/* [1] */
@@ -6397,19 +6437,19 @@
/* [14] */
/* fn atomic_i_increment[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, T, writable>, U, U) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(193),
+ /* overloads */ OverloadIndex(195),
},
{
/* [15] */
/* fn atomic_i_decrement[T : iu32, U : u32, S : workgroup_or_storage](ptr<S, T, writable>, U, U) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(193),
+ /* overloads */ OverloadIndex(195),
},
{
/* [16] */
/* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(194),
+ /* overloads */ OverloadIndex(196),
},
{
/* [17] */
@@ -6607,19 +6647,19 @@
/* [28] */
/* 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(195),
+ /* overloads */ OverloadIndex(197),
},
{
/* [29] */
/* 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(196),
+ /* overloads */ OverloadIndex(198),
},
{
/* [30] */
/* 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(197),
+ /* overloads */ OverloadIndex(199),
},
{
/* [31] */
@@ -6647,13 +6687,13 @@
/* [33] */
/* 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(198),
+ /* overloads */ OverloadIndex(200),
},
{
/* [34] */
/* fn vector_times_scalar[T : f32_f16, N : num](vec<N, T>, T) -> vec<N, T> */
/* num overloads */ 1,
- /* overloads */ OverloadIndex(199),
+ /* overloads */ OverloadIndex(201),
},
{
/* [35] */
@@ -6987,34 +7027,41 @@
},
{
/* [82] */
- /* fn s_dot(u32, u32, u32) -> i32 */
- /* num overloads */ 1,
- /* overloads */ OverloadIndex(200),
+ /* fn not<R : iu32>[A : iu32](A) -> R */
+ /* fn not<R : iu32>[A : iu32, N : num](vec<N, A>) -> vec<N, R> */
+ /* num overloads */ 2,
+ /* overloads */ OverloadIndex(192),
},
{
/* [83] */
- /* fn u_dot(u32, u32, u32) -> u32 */
- /* num overloads */ 1,
- /* overloads */ OverloadIndex(201),
- },
- {
- /* [84] */
- /* 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 */
+ /* fn s_dot(u32, u32, u32) -> i32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(202),
},
{
- /* [85] */
- /* 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) */
+ /* [84] */
+ /* fn u_dot(u32, u32, u32) -> u32 */
/* num overloads */ 1,
/* overloads */ OverloadIndex(203),
},
{
- /* [86] */
- /* 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> */
+ /* [85] */
+ /* 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] */
+ /* 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] */
+ /* 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),
+ },
};
// clang-format on
diff --git a/src/tint/lang/spirv/reader/lower/builtins.cc b/src/tint/lang/spirv/reader/lower/builtins.cc
index 0924c49..f267065 100644
--- a/src/tint/lang/spirv/reader/lower/builtins.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins.cc
@@ -203,12 +203,28 @@
case spirv::BuiltinFn::kShiftRightArithmetic:
ShiftRightArithmetic(builtin);
break;
+ case spirv::BuiltinFn::kNot:
+ Not(builtin);
+ break;
default:
TINT_UNREACHABLE() << "unknown spirv builtin: " << builtin->Func();
}
}
}
+ void Not(spirv::ir::BuiltinCall* call) {
+ auto* val = call->Args()[0];
+ auto* result_ty = call->Result(0)->Type();
+ b.InsertBefore(call, [&] {
+ auto* complement = b.Complement(val->Type(), val)->Result(0);
+ if (val->Type() != result_ty) {
+ complement = b.Bitcast(result_ty, complement)->Result(0);
+ }
+ call->Result(0)->ReplaceAllUsesWith(complement);
+ });
+ call->Destroy();
+ }
+
void ConvertSToF(spirv::ir::BuiltinCall* call) {
b.InsertBefore(call, [&] {
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 dd9ad92..3494df5 100644
--- a/src/tint/lang/spirv/reader/lower/builtins_test.cc
+++ b/src/tint/lang/spirv/reader/lower/builtins_test.cc
@@ -8501,5 +8501,233 @@
EXPECT_EQ(expect, str());
}
+TEST_F(SpirvReader_BuiltinsTest, Not_Scalar_Signed_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kNot, Vector{ty.i32()},
+ 1_i);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.not<i32> 1i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = complement 1i
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, Not_Scalar_Signed_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.u32(), spirv::BuiltinFn::kNot, Vector{ty.u32()},
+ 1_i);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.not<u32> 1i
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = complement 1i
+ %3:u32 = bitcast %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, Not_Scalar_Unsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.i32(), spirv::BuiltinFn::kNot, Vector{ty.i32()},
+ 8_u);
+ b.Return(ep);
+ });
+
+ auto src = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.not<i32> 8u
+ ret
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+ Run(Builtins);
+
+ auto expect = R"(
+%foo = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = complement 8u
+ %3:i32 = bitcast %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, Not_Vector_Signed_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kNot,
+ 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.not<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> = complement vec2<i32>(1i)
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, Not_Vector_Signed_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kNot,
+ 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.not<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> = complement vec2<i32>(1i)
+ %3:vec2<u32> = bitcast %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, Not_Vector_Unsigned_Signed) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<i32>(), spirv::BuiltinFn::kNot,
+ 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.not<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<u32> = complement vec2<u32>(8u)
+ %3:vec2<i32> = bitcast %2
+ ret
+ }
+}
+)";
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(SpirvReader_BuiltinsTest, Not_Vector_Unsigned_Unsigned) {
+ auto* ep = b.ComputeFunction("foo");
+
+ b.Append(ep->Block(), [&] { //
+ b.CallExplicit<spirv::ir::BuiltinCall>(ty.vec2<u32>(), spirv::BuiltinFn::kNot,
+ 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.not<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<u32> = complement vec2<u32>(8u)
+ 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 bf19d3f..f0d3cc9 100644
--- a/src/tint/lang/spirv/reader/parser/bit_test.cc
+++ b/src/tint/lang/spirv/reader/parser/bit_test.cc
@@ -1342,5 +1342,261 @@
)");
}
+TEST_F(SpirvParserTest, Not_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
+ %v2int = OpTypeVector %int 2
+ %v2uint = OpTypeVector %uint 2
+ %one = OpConstant %int 1
+ %eight = OpConstant %uint 8
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpNot %int %one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.not<i32> 1i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Not_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
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %eight = OpConstant %uint 8
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpNot %uint %one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.not<u32> 1i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Not_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
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %eight = OpConstant %uint 8
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpNot %int %eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.not<i32> 8u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Not_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
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %eight = OpConstant %uint 8
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpNot %uint %eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:u32 = spirv.not<u32> 8u
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Not_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
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %eight = OpConstant %uint 8
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpNot %v2int %v2one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.not<i32> vec2<i32>(1i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Not_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
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %eight = OpConstant %uint 8
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpNot %v2uint %v2one
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.not<u32> vec2<i32>(1i)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Not_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
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %eight = OpConstant %uint 8
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpNot %v2int %v2eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<i32> = spirv.not<i32> vec2<u32>(8u)
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Not_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
+ %v2uint = OpTypeVector %uint 2
+ %v2int = OpTypeVector %int 2
+ %one = OpConstant %int 1
+ %eight = OpConstant %uint 8
+ %v2one = OpConstantComposite %v2int %one %one
+ %v2eight = OpConstantComposite %v2uint %eight %eight
+ %void_fn = OpTypeFunction %void
+ %main = OpFunction %void None %void_fn
+ %main_start = OpLabel
+ %1 = OpNot %v2uint %v2eight
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:vec2<u32> = spirv.not<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 bbd6c3e..ef86fc9 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -958,6 +958,9 @@
case spv::Op::OpLogicalNot:
EmitUnary(inst, core::UnaryOp::kNot);
break;
+ case spv::Op::OpNot:
+ EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kNot);
+ break;
case spv::Op::OpShiftLeftLogical:
EmitSpirvExplicitBuiltinCall(inst, spirv::BuiltinFn::kShiftLeftLogical);
break;
diff --git a/src/tint/lang/spirv/spirv.def b/src/tint/lang/spirv/spirv.def
index ce0c4c44..3feb5fa 100644
--- a/src/tint/lang/spirv/spirv.def
+++ b/src/tint/lang/spirv/spirv.def
@@ -477,6 +477,9 @@
implicit(A: iu32, B: iu32) fn shift_right_arithmetic<R: iu32>(A, B) -> R
implicit(A: iu32, B: iu32, N: num) fn shift_right_arithmetic<R: iu32>(vec<N, A>, vec<N, B>) -> vec<N, R>
+implicit(A: iu32) fn not<R: iu32>(A) -> R
+implicit(A: iu32, N: num) fn not<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 2112311..eb55006 100644
--- a/src/tint/lang/spirv/writer/printer/printer.cc
+++ b/src/tint/lang/spirv/writer/printer/printer.cc
@@ -1587,6 +1587,9 @@
case BuiltinFn::kShiftRightArithmetic:
op = spv::Op::OpShiftRightArithmetic;
break;
+ case BuiltinFn::kNot:
+ op = spv::Op::OpNot;
+ break;
case spirv::BuiltinFn::kNone:
TINT_ICE() << "undefined spirv ir function";
}
