Remove the operators tanh since it was already added.

src/targets/gpu/CMakeLists.txt

@@ -15,7 +15,6 @@ add_library(migraphx_device
     device/sin.cpp
     device/sinh.cpp
     device/cosh.cpp
-    device/tanh.cpp
     device/add_relu.cpp
     device/contiguous.cpp
     device/mul.cpp
@@ -45,7 +44,6 @@ add_library(migraphx_gpu
     sin.cpp
     sinh.cpp
     cosh.cpp
-    tanh.cpp
     mul.cpp
     batchnorm.cpp
     write_literals.cpp

src/targets/gpu/device/tanh.cpp

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

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

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

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

-#ifndef MIGRAPH_GUARD_RTGLIB_TANH_HPP
-#define MIGRAPH_GUARD_RTGLIB_TANH_HPP
-
-#include <migraphx/gpu/lowering.hpp>
-#include <migraphx/manage_ptr.hpp>
-#include <migraphx/instruction.hpp>
-#include <migraphx/operators.hpp>
-#include <migraphx/generate.hpp>
-#include <migraphx/shape_for_each.hpp>
-#include <migraphx/gpu/miopen.hpp>
-#include <migraphx/gpu/hip.hpp>
-#include <migraphx/dfor.hpp>
-#include <migraphx/gpu/device/contiguous.hpp>
-#include <migraphx/gpu/device/tanh.hpp>
-#include <migraphx/iterator_for.hpp>
-#include <migraphx/gpu/rocblas.hpp>
-#include <migraphx/gpu/context.hpp>
-#include <migraphx/config.hpp>
-#include <utility>
-
-namespace migraphx {
-inline namespace MIGRAPH_INLINE_NS {
-namespace gpu {
-
-struct hip_tanh
-{
-    std::string name() const { return "gpu::tanh"; }
-    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
-} // namespace MIGRAPH_INLINE_NS
-} // namespace migraphx
-
-#endif

src/targets/gpu/lowering.cpp

@@ -22,7 +22,6 @@
 #include <migraphx/gpu/sin.hpp>
 #include <migraphx/gpu/sinh.hpp>
 #include <migraphx/gpu/cosh.hpp>
-#include <migraphx/gpu/tanh.hpp>
 #include <migraphx/gpu/mul.hpp>
 #include <migraphx/gpu/batchnorm.hpp>
 #include <migraphx/gpu/pooling.hpp>
@@ -83,10 +82,6 @@ struct miopen_apply
             {
                 check_shape(s, apply_cosh(it));
             }
-            else if(it->name() == "tanh")
-            {
-                check_shape(s, apply_tanh(it));
-            }
             else if(it->name() == "mul")
             {
                 check_shape(s, apply_mul(it));
@@ -203,12 +198,6 @@ struct miopen_apply
         return prog->replace_instruction(ins, hip_cosh{}, ins->inputs().at(0), output);
     }
 
-    instruction_ref apply_tanh(instruction_ref ins)
-    {
-        auto output = insert_allocation(ins, ins->get_shape());
-        return prog->replace_instruction(ins, hip_tanh{}, ins->inputs().at(0), output);
-    }
-
     instruction_ref apply_mul(instruction_ref ins)
     {
         auto output = insert_allocation(ins, ins->get_shape());

src/targets/gpu/tanh.cpp

-#include <migraphx/gpu/tanh.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 MIGRAPH_INLINE_NS {
-namespace gpu {
-
-shape hip_tanh::compute_shape(const std::vector<shape>& inputs) const
-{
-    check_shapes{inputs, *this}.has(2);
-    return inputs.at(0);
-}
-
-argument hip_tanh::compute(context& ctx, const shape&, const std::vector<argument>& args) const
-{
-    device::tanh(ctx.get_stream().get(), args[1], args[0]);
-    return args[1];
-}
-
-} // namespace gpu
-} // namespace MIGRAPH_INLINE_NS
-} // namespace migraphx

test/gpu/miopen.cpp

@@ -239,18 +239,6 @@ struct test_cosh
     }
 };
 
-struct test_tanh
-{
-    migraphx::program create_program() const
-    {
-        migraphx::program p;
-        migraphx::shape s{migraphx::shape::double_type, {16}};
-        auto x = p.add_parameter("x", s);
-        p.add_instruction(migraphx::op::tanh{}, x);
-        return p;
-    }
-};
-
 struct test_scale
 {
     migraphx::program create_program() const
@@ -894,7 +882,6 @@ int main()
     verify_program<test_sin>();
     verify_program<test_sinh>();
     verify_program<test_cosh>();
-    verify_program<test_tanh>();
     verify_program<test_scale>();
     verify_program<test_triadd>();
     verify_program<test_triadd2>();