Merged log and exp operators to current branch..

src/targets/gpu/exp.cpp

-#include <migraphx/gpu/exp.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/config.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-shape hip_exp::compute_shape(const std::vector<shape>& inputs) const
-{
-    check_shapes{inputs, *this}.has(2);
-    return inputs.at(0);
-}
-
-argument hip_exp::compute(context& ctx, const shape&, const std::vector<argument>& args) const
-{
-    device::exp(ctx.get_stream().get(), args[1], args[0]);
-    return args[1];
-}
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_EXP_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_exp
+struct hip_exp : unary_device<hip_exp, device::exp>
 {
-    std::string name() const { return "gpu::exp"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

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

@@ -2,6 +2,7 @@
 #define MIGRAPHX_GUARD_RTGLIB_LOG_HPP
 
 #include <migraphx/gpu/lowering.hpp>
+#include <migraphx/gpu/unary.hpp>
 #include <migraphx/manage_ptr.hpp>
 #include <migraphx/instruction.hpp>
 #include <migraphx/operators.hpp>
@@ -22,12 +23,8 @@ namespace migraphx {
 inline namespace MIGRAPHX_INLINE_NS {
 namespace gpu {
 
-struct hip_log
+struct hip_log : unary_device<hip_log, device::log>
 {
-    std::string name() const { return "gpu::log"; }
-    shape compute_shape(const std::vector<shape>& inputs) const;
-    argument compute(context&, const shape&, const std::vector<argument>& args) const;
-    int output_alias(const std::vector<shape>& shapes) const { return shapes.size() - 1; }
 };
 
 } // namespace gpu

src/targets/gpu/log.cpp

-#include <migraphx/gpu/log.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/config.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <utility>
-
-namespace migraphx {
-inline namespace MIGRAPHX_INLINE_NS {
-namespace gpu {
-
-shape hip_log::compute_shape(const std::vector<shape>& inputs) const
-{
-    check_shapes{inputs, *this}.has(2);
-    return inputs.at(0);
-}
-
-argument hip_log::compute(context& ctx, const shape&, const std::vector<argument>& args) const
-{
-    device::log(ctx.get_stream().get(), args[1], args[0]);
-    return args[1];
-}
-
-} // namespace gpu
-} // namespace MIGRAPHX_INLINE_NS
-} // namespace migraphx