added gpu version of max and min

src/targets/gpu/CMakeLists.txt

@@ -17,6 +17,8 @@ add_library(migraphx_device
     device/contiguous.cpp
     device/mul.cpp
     device/concat.cpp
+    device/max.cpp
+    device/min.cpp
 )
 set_target_properties(migraphx_device PROPERTIES EXPORT_NAME device)
 rocm_clang_tidy_check(migraphx_device)
@@ -48,6 +50,8 @@ add_library(migraphx_gpu
     tanh.cpp
     abs.cpp
     elu.cpp
+    max.cpp
+    min.cpp
 )
 set_target_properties(migraphx_gpu PROPERTIES EXPORT_NAME gpu)
 rocm_clang_tidy_check(migraphx_gpu)

src/targets/gpu/device/max.cpp

+#include <migraphx/gpu/device/max.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 max(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2)
+{
+    nary(stream, result, arg1, arg2)([](auto x, auto y) { return std::max(to_hip_type(x), to_hip_type(y)); });
+}
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx

src/targets/gpu/device/min.cpp

+#include <migraphx/gpu/device/min.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 min(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2)
+{
+    nary(stream, result, arg1, arg2)([](auto x, auto y) { return std::min(to_hip_type(x), to_hip_type(y)); });
+}
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx

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

+
+#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_MAX_HPP
+#define MIGRAPH_GUARD_RTGLIB_DEVICE_MAX_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 max(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2);
+
+void max(hipStream_t stream,
+         const argument& result,
+         const argument& arg1,
+         const argument& arg2,
+         const argument& arg3);
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx
+
+#endif

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

+
+#ifndef MIGRAPH_GUARD_RTGLIB_DEVICE_MIN_HPP
+#define MIGRAPH_GUARD_RTGLIB_DEVICE_MIN_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 min(hipStream_t stream, const argument& result, const argument& arg1, const argument& arg2);
+
+void min(hipStream_t stream,
+         const argument& result,
+         const argument& arg1,
+         const argument& arg2,
+         const argument& arg3);
+
+} // namespace device
+} // namespace gpu
+} // namespace MIGRAPH_INLINE_NS
+} // namespace migraphx
+
+#endif

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

+#ifndef MIGRAPH_GUARD_RTGLIB_MAX_HPP
+#define MIGRAPH_GUARD_RTGLIB_MAX_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/config.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/max.hpp>
+#include <migraphx/iterator_for.hpp>
+#include <migraphx/gpu/rocblas.hpp>
+#include <migraphx/gpu/context.hpp>
+#include <utility>
+
+namespace migraphx {
+inline namespace MIGRAPH_INLINE_NS {
+namespace gpu {
+
+struct hip_max
+{
+    std::string name() const { return "gpu::max"; }
+    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/include/migraphx/gpu/min.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_MIN_HPP
+#define MIGRAPH_GUARD_RTGLIB_MIN_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/config.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/min.hpp>
+#include <migraphx/iterator_for.hpp>
+#include <migraphx/gpu/rocblas.hpp>
+#include <migraphx/gpu/context.hpp>
+#include <utility>
+
+namespace migraphx {
+inline namespace MIGRAPH_INLINE_NS {
+namespace gpu {
+
+struct hip_min
+{
+    std::string name() const { return "gpu::min"; }
+    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

@@ -25,6 +25,8 @@
 #include <migraphx/gpu/add.hpp>
 #include <migraphx/gpu/sin.hpp>
 #include <migraphx/gpu/mul.hpp>
+#include <migraphx/gpu/max.hpp>
+#include <migraphx/gpu/min.hpp>
 #include <migraphx/gpu/batchnorm.hpp>
 #include <migraphx/gpu/pooling.hpp>
 #include <migraphx/gpu/gemm.hpp>
@@ -63,6 +65,8 @@ struct miopen_apply
         apply_map["add"]                  = &miopen_apply::apply_add;
         apply_map["sin"]                  = &miopen_apply::apply_sin;
         apply_map["mul"]                  = &miopen_apply::apply_mul;
+        apply_map["max"]                  = &miopen_apply::apply_max;
+        apply_map["min"]                  = &miopen_apply::apply_min;
         apply_map["dot"]                  = &miopen_apply::apply_dot;
         apply_map["contiguous"]           = &miopen_apply::apply_contiguous;
         apply_map["concat"]               = &miopen_apply::apply_concat;
@@ -204,6 +208,20 @@ struct miopen_apply
             ins, hip_mul{}, ins->inputs().at(0), ins->inputs().at(1), output);
     }
 
+    instruction_ref apply_max(instruction_ref ins)
+    {
+        auto output = insert_allocation(ins, ins->get_shape());
+        return prog->replace_instruction(
+            ins, hip_max{}, ins->inputs().at(0), ins->inputs().at(1), output);
+    }
+
+    instruction_ref apply_min(instruction_ref ins)
+    {
+        auto output = insert_allocation(ins, ins->get_shape());
+        return prog->replace_instruction(
+            ins, hip_min{}, ins->inputs().at(0), ins->inputs().at(1), output);
+    }
+
     instruction_ref apply_dot(instruction_ref ins)
     {
         auto&& op   = any_cast<op::dot>(ins->get_operator());

src/targets/gpu/max.cpp

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

src/targets/gpu/min.cpp

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