Ceil floor operators (#375)

* add two operators ceil and floor

* clang format

* add unit test for the ceil and floor operators

* remove unintended code

src/include/migraphx/op/ceil.hpp

+#ifndef MIGRAPHX_GUARD_OPERATORS_CEIL_HPP
+#define MIGRAPHX_GUARD_OPERATORS_CEIL_HPP
+
+#include <migraphx/op/unary.hpp>
+#include <cmath>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace op {
+
+struct ceil : unary<ceil>
+{
+    auto apply() const
+    {
+        return [](auto x) { return std::ceil(x); };
+    }
+};
+
+} // namespace op
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/include/migraphx/op/floor.hpp

+#ifndef MIGRAPHX_GUARD_OPERATORS_FLOOR_HPP
+#define MIGRAPHX_GUARD_OPERATORS_FLOOR_HPP
+
+#include <migraphx/op/unary.hpp>
+#include <cmath>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace op {
+
+struct floor : unary<floor>
+{
+    auto apply() const
+    {
+        return [](auto x) { return std::floor(x); };
+    }
+};
+
+} // namespace op
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/include/migraphx/operators.hpp

@@ -14,6 +14,7 @@
 #include <migraphx/op/binary.hpp>
 #include <migraphx/op/broadcast.hpp>
 #include <migraphx/op/capture.hpp>
+#include <migraphx/op/ceil.hpp>
 #include <migraphx/op/clip.hpp>
 #include <migraphx/op/common.hpp>
 #include <migraphx/op/concat.hpp>
@@ -28,6 +29,7 @@
 #include <migraphx/op/erf.hpp>
 #include <migraphx/op/exp.hpp>
 #include <migraphx/op/flatten.hpp>
+#include <migraphx/op/floor.hpp>
 #include <migraphx/op/gather.hpp>
 #include <migraphx/op/gru.hpp>
 #include <migraphx/op/identity.hpp>

src/onnx/onnx.cpp

@@ -57,6 +57,8 @@ struct onnx_parser
         add_generic_op("Sqrt", op::sqrt{});
         add_generic_op("Round", op::round{});
         add_generic_op("Sign", op::sign{});
+        add_generic_op("Ceil", op::ceil{});
+        add_generic_op("Floor", op::floor{});
 
         add_binary_op("Add", op::add{});
         add_binary_op("Div", op::div{});

src/targets/gpu/CMakeLists.txt

@@ -54,6 +54,8 @@ add_library(migraphx_device
     device/pow.cpp
     device/sqdiff.cpp
     device/sign.cpp
+    device/ceil.cpp
+    device/floor.cpp
 )
 set_target_properties(migraphx_device PROPERTIES EXPORT_NAME device)
 rocm_set_soversion(migraphx_device ${PROJECT_VERSION})

src/targets/gpu/device/ceil.cpp

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

src/targets/gpu/device/floor.cpp

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

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

+#ifndef MIGRAPHX_GUARD_RTGLIB_CEIL_HPP
+#define MIGRAPHX_GUARD_RTGLIB_CEIL_HPP
+
+#include <migraphx/gpu/oper.hpp>
+#include <migraphx/gpu/device/ceil.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+struct hip_ceil : unary_device<hip_ceil, device::ceil>
+{
+};
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

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

+#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_CEIL_HPP
+#define MIGRAPHX_GUARD_RTGLIB_DEVICE_CEIL_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 ceil(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/device/floor.hpp

+#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_FLOOR_HPP
+#define MIGRAPHX_GUARD_RTGLIB_DEVICE_FLOOR_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 floor(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/floor.hpp

+#ifndef MIGRAPHX_GUARD_RTGLIB_FLOOR_HPP
+#define MIGRAPHX_GUARD_RTGLIB_FLOOR_HPP
+
+#include <migraphx/gpu/oper.hpp>
+#include <migraphx/gpu/device/floor.hpp>
+
+namespace migraphx {
+inline namespace MIGRAPHX_INLINE_NS {
+namespace gpu {
+
+struct hip_floor : unary_device<hip_floor, device::floor>
+{
+};
+
+} // namespace gpu
+} // namespace MIGRAPHX_INLINE_NS
+} // namespace migraphx
+
+#endif

src/targets/gpu/lowering.cpp

@@ -55,6 +55,8 @@
 #include <migraphx/gpu/clip.hpp>
 #include <migraphx/gpu/reduce_sum.hpp>
 #include <migraphx/gpu/round.hpp>
+#include <migraphx/gpu/ceil.hpp>
+#include <migraphx/gpu/floor.hpp>
 #include <migraphx/gpu/rsqrt.hpp>
 #include <migraphx/gpu/sqrt.hpp>
 #include <migraphx/gpu/reduce_mean.hpp>
@@ -120,6 +122,8 @@ struct miopen_apply
         add_generic_op<hip_relu>("relu");
         add_generic_op<hip_sign>("sign");
         add_generic_op<hip_sigmoid>("sigmoid");
+        add_generic_op<hip_ceil>("ceil");
+        add_generic_op<hip_floor>("floor");
 
         add_extend_op<miopen_contiguous, op::contiguous>("contiguous");
         add_extend_op<hip_concat, op::concat>("concat");

test/cpu_ops_test.cpp

@@ -984,7 +984,6 @@ TEST_CASE(maxpool_test)
     p.add_instruction(migraphx::op::pooling{"max", {{0, 0}}, {{2, 2}}, {{3, 2}}}, al);
     p.compile(migraphx::cpu::target{});
     auto result = p.eval({});
-    // std::cout << result.get_shape() << std::endl;
     std::vector<float> results_vector(36);
     result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
     EXPECT(migraphx::verify_range(results_vector, c));
@@ -2129,15 +2128,38 @@ TEST_CASE(round_test)
     auto result = p.eval({});
     std::vector<float> results_vector;
     result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
-    for(auto v : results_vector)
-    {
-        std::cout << v << "\t";
-    }
-    std::cout << std::endl;
     std::vector<float> gold = {1.0, 2.0, 2.0, -1.0, -2.0, -2.0, 0.0, 2.0, -2.0};
     EXPECT(migraphx::verify_range(results_vector, gold));
 }
 
+TEST_CASE(ceil_test)
+{
+    migraphx::program p;
+    migraphx::shape s{migraphx::shape::float_type, {9}};
+    auto l = p.add_literal(migraphx::literal{s, {1.1, 1.5, 1.6, -1.1, -1.5, -1.6, 0.0, 2.0, -2.0}});
+    p.add_instruction(migraphx::op::ceil{}, l);
+    p.compile(migraphx::cpu::target{});
+    auto result = p.eval({});
+    std::vector<float> results_vector;
+    result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+    std::vector<float> gold = {2.0, 2.0, 2.0, -1.0, -1.0, -1.0, 0.0, 2.0, -2.0};
+    EXPECT(migraphx::verify_range(results_vector, gold));
+}
+
+TEST_CASE(floor_test)
+{
+    migraphx::program p;
+    migraphx::shape s{migraphx::shape::float_type, {9}};
+    auto l = p.add_literal(migraphx::literal{s, {1.1, 1.5, 0.6, -1.1, -1.5, -0.6, 0.0, 2.0, -2.0}});
+    p.add_instruction(migraphx::op::floor{}, l);
+    p.compile(migraphx::cpu::target{});
+    auto result = p.eval({});
+    std::vector<float> results_vector;
+    result.visit([&](auto output) { results_vector.assign(output.begin(), output.end()); });
+    std::vector<float> gold = {1.0, 1.0, 0.0, -2.0, -2.0, -1.0, -0.0, 2.0, -2.0};
+    EXPECT(migraphx::verify_range(results_vector, gold));
+}
+
 TEST_CASE(op_capture)
 {
     migraphx::program p;

test/gpu/ops_test.cpp

@@ -3911,6 +3911,32 @@ struct test_round : verify_program<test_round>
     };
 };
 
+struct test_ceil : verify_program<test_ceil>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+
+        migraphx::shape s{migraphx::shape::double_type, {2, 3, 4, 6}};
+        auto param = p.add_parameter("x", s);
+        p.add_instruction(migraphx::op::ceil{}, param);
+        return p;
+    };
+};
+
+struct test_floor : verify_program<test_floor>
+{
+    migraphx::program create_program() const
+    {
+        migraphx::program p;
+
+        migraphx::shape s{migraphx::shape::float_type, {2, 3, 4, 6}};
+        auto param = p.add_parameter("x", s);
+        p.add_instruction(migraphx::op::floor{}, param);
+        return p;
+    };
+};
+
 struct test_convert : verify_program<test_convert>
 {
     migraphx::program create_program() const

test/onnx/ceil_test.onnx

+	ceil_test:;
+
+
x
y"Ceil	ceil_testZ
+
x
+
+

+
+b
+
y
+
+

+
+B	
\ No newline at end of file

test/onnx/floor_test.onnx

+
+floor_test:=
+
+
x
y"Floor
+floor_testZ
+
x
+
+

+
+b
+
y
+
+

+
+B	
\ No newline at end of file

test/onnx/gen_onnx.py

@@ -153,6 +153,18 @@ def cast_test():
 
     return ([node], [x], [y])
 
+@onnx_test
+def ceil_test():
+    x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [10])
+    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [10])
+
+    node = onnx.helper.make_node(
+        'Ceil',
+        inputs=['x'],
+        outputs=['y'],
+    )
+
+    return ([node], [x], [y])
 
 @onnx_test
 def clip_test():
@@ -635,6 +647,18 @@ def flatten_test():
 
     return ([node, node2], [x], [y, y2])
 
+@onnx_test
+def floor_test():
+    x = helper.make_tensor_value_info('x', TensorProto.FLOAT, [10])
+    y = helper.make_tensor_value_info('y', TensorProto.FLOAT, [10])
+
+    node = onnx.helper.make_node(
+        'Floor',
+        inputs=['x'],
+        outputs=['y'],
+    )
+
+    return ([node], [x], [y])
 
 @onnx_test
 def gather_test():

test/onnx/onnx_test.cpp

@@ -112,6 +112,17 @@ TEST_CASE(cast_test)
     EXPECT(p == prog);
 }
 
+TEST_CASE(ceil_test)
+{
+    migraphx::program p;
+    auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {10}});
+    p.add_instruction(migraphx::op::ceil{}, input);
+
+    auto prog = migraphx::parse_onnx("ceil_test.onnx");
+
+    EXPECT(p == prog);
+}
+
 TEST_CASE(clip_test)
 {
     migraphx::program p;
@@ -402,6 +413,17 @@ TEST_CASE(flatten_test)
     EXPECT(p == prog);
 }
 
+TEST_CASE(floor_test)
+{
+    migraphx::program p;
+    auto input = p.add_parameter("x", migraphx::shape{migraphx::shape::float_type, {10}});
+    p.add_instruction(migraphx::op::floor{}, input);
+
+    auto prog = migraphx::parse_onnx("floor_test.onnx");
+
+    EXPECT(p == prog);
+}
+
 TEST_CASE(gather_test)
 {
     migraphx::program p;