changes for not operator (#735)

* changes for not operator

* changed name of the op from unary_not to not

* Added tests for op and onnx parsing

* reordering not_test in onnx_test.cpp

* not operator -- gpu implementation

* added bool test for not operator

* Added test and missing links for not operator on GPU

* typo fix

* adding .onnx test files for not operator

* formatting
Co-authored-by: Shucai Xiao <shucai@gmail.com>
Co-authored-by: mvermeulen <5479696+mvermeulen@users.noreply.github.com>

src/CMakeLists.txt

@@ -145,6 +145,7 @@ register_migraphx_ops(
     tanh
     tan
     transpose
+    unary_not
     undefined
     unknown
     unsqueeze

src/include/migraphx/op/unary_not.hpp

+#ifndef MIGRAPHX_GUARD_OPERATORS_UNARY_NOT_HPP
+#define MIGRAPHX_GUARD_OPERATORS_UNARY_NOT_HPP
+
+#include <migraphx/op/unary.hpp>
+#include <migraphx/operation.hpp>
+#include <migraphx/check_shapes.hpp>
+#include <migraphx/config.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace op {
+
+struct unary_not : unary<unary_not>
+{
+    auto apply() const
+    {
+        return [](auto x) { return not x; };
+    }
+
+    std::string name() const { return "not"; }
+};
+
+} // namespace op
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/include/migraphx/operators.hpp

@@ -91,6 +91,7 @@
 #include <migraphx/op/tan.hpp>
 #include <migraphx/op/transpose.hpp>
 #include <migraphx/op/unary.hpp>
+#include <migraphx/op/unary_not.hpp>
 #include <migraphx/op/undefined.hpp>
 #include <migraphx/op/unknown.hpp>
 #include <migraphx/op/unsqueeze.hpp>

src/onnx/parse_generic_op.cpp

@@ -45,6 +45,7 @@ struct parse_generic_op : op_parser<parse_generic_op>
                 {"Squeeze", "squeeze"},
                 {"Tan", "tan"},
                 {"Tanh", "tanh"},
+                {"Not", "not"},
                 {"Unsqueeze", "unsqueeze"}};
     }
 

src/targets/gpu/CMakeLists.txt

@@ -74,6 +74,7 @@ add_library(migraphx_device
     device/tan.cpp
     device/tanh.cpp
     device/rnn_variable_seq_lens.cpp
+    device/unary_not.cpp
     device/equal.cpp
 )
 set_target_properties(migraphx_device PROPERTIES EXPORT_NAME device)
@@ -199,6 +200,7 @@ register_migraphx_gpu_ops(hip_
     sub
     tanh
     tan
+    unary_not
 )
 register_migraphx_gpu_ops(miopen_
     abs

src/targets/gpu/device/unary_not.cpp

+#include <migraphx/gpu/device/unary_not.hpp>
+#include <migraphx/gpu/device/nary.hpp>
+#include <migraphx/type_traits.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+namespace device {
+
+void unary_not(hipStream_t stream,
+                 const argument& result,
+                 const argument& arg
+                 )
+{
+    nary(stream, result, arg)(
+        [](auto x) __device__ { return not x; });
+}
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx

src/targets/gpu/include/migraphx/gpu/device/unary_not.hpp

+#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_UNARY_NOT_HPP
+#define MIGRAPHX_GUARD_RTGLIB_DEVICE_UNARY_NOT_HPP
+
+#include <migraphx/argument.hpp>
+#include <migraphx/config.hpp>
+#include <hip/hip_runtime_api.h>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+namespace device {
+
+void unary_not(hipStream_t stream,
+                 const argument& result,
+                 const argument& arg);
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/targets/gpu/include/migraphx/gpu/unary_not.hpp

+#ifndef MIGRAPHX_GUARD_RTGLIB_UNARY_NOT_HPP
+#define MIGRAPHX_GUARD_RTGLIB_UNARY_NOT_HPP
+
+#include <migraphx/gpu/oper.hpp>
+#include <migraphx/gpu/device/unary_not.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+struct hip_unary_not : unary_device<hip_unary_not, device::unary_not>
+{
+    std::string name() const { return "gpu::not"; }
+};
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/targets/gpu/lowering.cpp

@@ -36,6 +36,7 @@
 #include <migraphx/gpu/miopen.hpp>
 #include <migraphx/gpu/quant_convolution.hpp>
 #include <migraphx/gpu/rocblas.hpp>
+#include <migraphx/gpu/unary_not.hpp>
 #include <migraphx/iterator_for.hpp>
 #include <migraphx/program.hpp>
 #include <utility>
@@ -121,6 +122,7 @@ struct miopen_apply
         add_generic_op("max");
         add_generic_op("min");
         add_generic_op("mul");
+        add_generic_op("not");
         add_generic_op("pow");
         add_generic_op("prelu");
         add_generic_op("recip");
@@ -137,6 +139,7 @@ struct miopen_apply
         add_generic_op("tan");
         add_generic_op("tanh");
 
+
         add_extend_op("abs");
         add_extend_op("argmax");
         add_extend_op("argmin");

test/onnx/gen_onnx.py

@@ -2143,6 +2143,26 @@ def neg_test():
     return ([node], [x], [y])
 
 
+@onnx_test
+def not_test():
+    x = helper.make_tensor_value_info('0', TensorProto.INT32, [4])
+    y = helper.make_tensor_value_info('1', TensorProto.INT32, [4])
+
+    node = onnx.helper.make_node('Not', inputs=['0'], outputs=['1'])
+
+    return ([node], [x], [y])
+
+
+@onnx_test
+def not_bool_test():
+    x = helper.make_tensor_value_info('0', TensorProto.BOOL, [4])
+    y = helper.make_tensor_value_info('1', TensorProto.BOOL, [4])
+
+    node = onnx.helper.make_node('Not', inputs=['0'], outputs=['1'])
+
+    return ([node], [x], [y])
+
+
 @onnx_test
 def no_pad_test():
     x = helper.make_tensor_value_info('0', TensorProto.FLOAT, [2, 2])

test/onnx/not_bool_test.onnx

+
not_bool_test:>
+
+
0
1"Not
not_bool_testZ
+
0
+
+	
+b
+
1
+
+	
+B
\ No newline at end of file

test/onnx/not_test.onnx

+not_test:9
+
+
0
1"Notnot_testZ
+
0
+
+
+b
+
1
+
+
+B
\ No newline at end of file

test/onnx/onnx_test.cpp

@@ -1684,6 +1684,7 @@ TEST_CASE(logical_or_test)
     EXPECT(p == prog);
 }
 
+
 TEST_CASE(logical_xor_bcast_test)
 {
     migraphx::program p;
@@ -1984,6 +1985,32 @@ TEST_CASE(nonzero_int_test)
     EXPECT(p == prog);
 }
 
+TEST_CASE(not_test)
+{
+    migraphx::program p;
+    auto* mm = p.get_main_module();
+    auto l0  = mm->add_parameter("0", migraphx::shape{migraphx::shape::int32_type, {4}});
+    auto ret = mm->add_instruction(migraphx::make_op("not"), l0);
+    mm->add_return({ret});
+
+    auto prog = migraphx::parse_onnx("not_test.onnx");
+
+    EXPECT(p == prog);
+}
+
+TEST_CASE(not_bool_test)
+{
+    migraphx::program p;
+    auto* mm = p.get_main_module();
+    auto l0  = mm->add_parameter("0", migraphx::shape{migraphx::shape::bool_type, {4}});
+    auto ret = mm->add_instruction(migraphx::make_op("not"), l0);
+    mm->add_return({ret});
+
+    auto prog = migraphx::parse_onnx("not_bool_test.onnx");
+
+    EXPECT(p == prog);
+}
+
 TEST_CASE(onehot_test)
 {
     migraphx::program p;

test/ref_ops_test.cpp

@@ -1838,6 +1838,40 @@ TEST_CASE(logical_and_test)
     EXPECT(migraphx::verify_range(results_vector, gold));
 }
 
+TEST_CASE(not_test)
+{
+    //int32
+    {
+        migraphx::program p;
+        auto* mm = p.get_main_module();
+        migraphx::shape s{migraphx::shape::int32_type, {4}};
+        auto l1 = mm->add_literal(migraphx::literal{s, {0, 8, 1, -32}});
+        mm->add_instruction(migraphx::make_op("not"), l1);
+        p.compile(migraphx::ref::target{});
+        auto result = p.eval({}).back();
+        std::vector<char> results_vector;
+        result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+        std::vector<char> gold = {1, 0, 0, 0};
+        EXPECT(migraphx::verify_range(results_vector, gold));
+    }
+
+    //bool
+    {
+        migraphx::program p;
+        auto* mm = p.get_main_module();
+        migraphx::shape s{migraphx::shape::bool_type, {4}};
+        auto l1 = mm->add_literal(migraphx::literal{s, {0, 0, 1, 1}});
+        mm->add_instruction(migraphx::make_op("not"), l1);
+        p.compile(migraphx::ref::target{});
+        auto result = p.eval({}).back();
+        std::vector<char> results_vector;
+        result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+        std::vector<char> gold = {1, 1, 0, 0};
+        EXPECT(migraphx::verify_range(results_vector, gold));
+    }
+
+}
+
 TEST_CASE(logical_or_test)
 {
     migraphx::program p;

test/verify/test_not.cpp

+
+#include "verify_program.hpp"
+#include <migraphx/program.hpp>
+#include <migraphx/generate.hpp>
+#include <migraphx/make_op.hpp>
+
+struct test_not : verify_program<test_not>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+        auto* mm = p.get_main_module();
+        migraphx::shape s{migraphx::shape::bool_type, {4}};
+        auto x = mm->add_parameter("x", s);
+        mm->add_instruction(migraphx::make_op("not"), x);
+        return p;
+    }
+};