feat: add support for the Neg operator (elementwise -x)

Author: pixelspark

README.md

@@ -291,7 +291,7 @@ fn test_matmul_square_matrix() {
 |<a href="https://github.com/onnx/onnx/blob/main/docs/Operators.md#Mod">Mod</a>|<a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Mod-13">13</a>, <a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Mod-10">10</a>|✅|✅|
 |<a href="https://github.com/onnx/onnx/blob/main/docs/Operators.md#Mul">Mul</a>|<a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Mul-14">14</a>, <a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Mul-13">13</a>, <a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Mul-7">7</a>, <a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Mul-6">6</a>, <a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Mul-1">1</a>|✅|✅|
 |<a href="https://github.com/onnx/onnx/blob/main/docs/Operators.md#Multinomial">Multinomial</a>|<a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Multinomial-7">7</a>|
-|<a href="https://github.com/onnx/onnx/blob/main/docs/Operators.md#Neg">Neg</a>|<a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Neg-13">13</a>, <a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Neg-6">6</a>, <a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Neg-1">1</a>|
+|<a href="https://github.com/onnx/onnx/blob/main/docs/Operators.md#Neg">Neg</a>|<a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Neg-13">13</a>, <a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Neg-6">6</a>, <a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Neg-1">1</a>|✅|✅|
 |<a href="https://github.com/onnx/onnx/blob/main/docs/Operators.md#NonMaxSuppression">NonMaxSuppression</a>|<a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#NonMaxSuppression-11">11</a>, <a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#NonMaxSuppression-10">10</a>|
 |<a href="https://github.com/onnx/onnx/blob/main/docs/Operators.md#NonZero">NonZero</a>|<a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#NonZero-13">13</a>, <a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#NonZero-9">9</a>|
 |<a href="https://github.com/onnx/onnx/blob/main/docs/Operators.md#Not">Not</a>|<a href="https://github.com/onnx/onnx/blob/main/docs/Changelog.md#Not-1">1</a>|✅|

wonnx-py/tests/test_onnx_backend.py

@@ -155,6 +155,8 @@ def do_enforce_test_coverage_safelist(model):  # type: (ModelProto) -> bool
 backend_test.include(f"test_sub_bcast_[a-z,_]*")
 backend_test.include(f"test_pow_bcast_[a-z,_]*")
 backend_test.include(f"test_transpose[a-z,_]*")
+backend_test.include(f"test_neg_[a-z,_]*")
+backend_test.include(f"test_reciprocal_[a-z,_]*")
 
 # Don't support 'bool' type
 # backend_test.include(f"test_and_bcast[a-z0-9,_]*")

wonnx/src/compiler.rs

@@ -288,7 +288,7 @@ pub fn compile(
         // Map simple function
         "Abs" | "Acos" | "Asin" | "Atan" | "Ceil" | "Cos" | "Cosh" | "Exp" | "Floor" | "Log"
         | "Round" | "Sign" | "Sin" | "Sinh" | "Sqrt" | "Tan" | "Tanh" | "Reciprocal" | "Acosh"
-        | "Asinh" | "Atanh" => {
+        | "Asinh" | "Atanh" | "Neg" => {
             let (x_threads, workgroup_size_x) = workgroup_size(
                 ceil(output_lengths[0], 4),
                 MAX_COMPUTE_WORKGROUPS_PER_DIMENSION,

wonnx/templates/endomorphism/map.wgsl

@@ -13,6 +13,10 @@ fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
 		let one_scalar = {{scalar_type}}(1);
 		let one = Vec4(one_scalar, one_scalar, one_scalar, one_scalar);
 		output_0.data[gidx] = one / (input_0.data[gidx]);
+	{% elif op_type == "Neg" %}
+		let zero_scalar = {{scalar_type}}(0);
+		let zeroes = Vec4(zero_scalar, zero_scalar, zero_scalar, zero_scalar);
+		output_0.data[gidx] = zeroes - (input_0.data[gidx]);
 	{% elif op_type == "Tanh" %}
 		{# Tanh will produce NaNs when fed with inputs that are much larger than +10.0 or smaller than -10.0. As the output
 		for these inputs converges to 1.0 and -1.0 respectively, we clamp the inputs first. #}