Skip to content

GitLab

"test/git@developer.sourcefind.cn:gaoqiong/migraphx.git" did not exist on "701c20147e6f4c5cf28aa3c8f2a6e0eae4d3ba77"
Unverified Commit fef8086c authored by mvermeulen's avatar mvermeulen Committed by GitHub
Browse files

Merge pull request #326 from ROCmSoftwarePlatform/rewrite_relu_tanh_device_func 

Rewrite relu tanh device to avoid calling MIOpen
parents 1fe84f2a ba45fce2
Hide whitespace changes
Inline Side-by-side
Showing with 85 additions and 116 deletions
src/targets/gpu/CMakeLists.txt
View file @ fef8086c
......@@ -24,9 +24,11 @@ add_library(migraphx_device
device/tan.cpp
device/sinh.cpp
device/cosh.cpp
device/tanh.cpp
device/asin.cpp
device/acos.cpp
device/atan.cpp
device/relu.cpp
device/add_relu.cpp
device/contiguous.cpp
device/logsoftmax.cpp
......@@ -68,9 +70,7 @@ add_library(migraphx_gpu
logsoftmax.cpp
contiguous.cpp
concat.cpp
relu.cpp
leaky_relu.cpp
tanh.cpp
batchnorm.cpp
write_literals.cpp
rocblas.cpp
......
src/targets/gpu/device/relu.cpp 0 → 100644
View file @ fef8086c
#include <migraphx/gpu/device/relu.hpp>
#include <migraphx/gpu/device/nary.hpp>
namespace migraphx {
inline namespace MIGRAPHX_INLINE_NS {
namespace gpu {
namespace device {
void relu(hipStream_t stream, const argument& result, const argument& arg)
{
nary(stream, result, arg)([](auto x) { return std::max<decltype(x)>(0, x); });
}
} // namespace device
} // namespace gpu
} // namespace MIGRAPHX_INLINE_NS
} // namespace migraphx
src/targets/gpu/device/tanh.cpp 0 → 100644
View file @ fef8086c
#include <migraphx/gpu/device/tanh.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 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 MIGRAPHX_INLINE_NS
} // namespace migraphx
src/targets/gpu/include/migraphx/gpu/device/relu.hpp 0 → 100644
View file @ fef8086c
#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_RELU_HPP
#define MIGRAPHX_GUARD_RTGLIB_DEVICE_RELU_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 relu(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/tanh.hpp 0 → 100644
View file @ fef8086c
#ifndef MIGRAPHX_GUARD_RTGLIB_DEVICE_TANH_HPP
#define MIGRAPHX_GUARD_RTGLIB_DEVICE_TANH_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 tanh(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/relu.hpp
View file @ fef8086c
#ifndef MIGRAPHX_GUARD_RTGLIB_RELU_HPP
#define MIGRAPHX_GUARD_RTGLIB_RELU_HPP
#include <migraphx/shape.hpp>
#include <migraphx/gpu/miopen.hpp>
#include <migraphx/gpu/oper.hpp>
#include <migraphx/gpu/device/relu.hpp>
namespace migraphx {
inline namespace MIGRAPHX_INLINE_NS {
......@@ -10,24 +10,8 @@ namespace gpu {
struct context;
struct miopen_relu
struct hip_relu : unary_device<hip_relu, device::relu>
{
shared<activation_descriptor> ad;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return gpu::reflect(self.ad.get(), f);
}
std::string name() const { return "gpu::relu"; }
shape compute_shape(const std::vector<shape>& inputs) const;
argument
compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
std::ptrdiff_t output_alias(const std::vector<shape>& shapes) const
{
return shapes.size() - 1;
}
};
} // namespace gpu
......
src/targets/gpu/include/migraphx/gpu/tanh.hpp
View file @ fef8086c
#ifndef MIGRAPHX_GUARD_RTGLIB_TANH_HPP
#define MIGRAPHX_GUARD_RTGLIB_TANH_HPP
#include <migraphx/shape.hpp>
#include <migraphx/gpu/miopen.hpp>
#include <migraphx/gpu/oper.hpp>
#include <migraphx/gpu/device/tanh.hpp>
namespace migraphx {
inline namespace MIGRAPHX_INLINE_NS {
namespace gpu {
struct context;
struct miopen_tanh
struct hip_tanh : unary_device<hip_tanh, device::tanh>
{
shared<activation_descriptor> ad;
template <class Self, class F>
static auto reflect(Self& self, F f)
{
return gpu::reflect(self.ad.get(), f);
}
std::string name() const { return "gpu::tanh"; }
shape compute_shape(const std::vector<shape>& inputs) const;
argument
compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
std::ptrdiff_t output_alias(const std::vector<shape>& shapes) const
{
return shapes.size() - 1;
}
};
} // namespace gpu
......
src/targets/gpu/lowering.cpp
View file @ fef8086c
......@@ -83,10 +83,8 @@ struct miopen_apply
void init()
{
this->last = instruction::get_output_alias(std::prev(prog->end()));
add_miopen_simple_op<miopen_relu>("relu", make_relu);
add_miopen_simple_op<miopen_sigmoid>("sigmoid", make_sigmoid);
add_miopen_simple_op<miopen_abs>("abs", make_abs);
add_miopen_simple_op<miopen_tanh>("tanh", make_tanh);
add_miopen_extend_op<miopen_leaky_relu, op::leaky_relu>("leaky_relu", make_leaky_relu);
add_miopen_extend_op<miopen_elu, op::elu>("elu", make_elu);
......@@ -101,6 +99,7 @@ struct miopen_apply
add_generic_op<hip_tan>("tan");
add_generic_op<hip_sinh>("sinh");
add_generic_op<hip_cosh>("cosh");
add_generic_op<hip_tanh>("tanh");
add_generic_op<hip_asin>("asin");
add_generic_op<hip_acos>("acos");
add_generic_op<hip_atan>("atan");
......@@ -112,6 +111,7 @@ struct miopen_apply
add_generic_op<hip_rsqrt>("rsqrt");
add_generic_op<hip_pow>("pow");
add_generic_op<hip_sqdiff>("sqdiff");
add_generic_op<hip_relu>("relu");
add_generic_op<hip_sign>("sign");
add_extend_op<miopen_gemm, op::dot>("dot");
......
src/targets/gpu/relu.cpp deleted 100644 → 0
View file @ 1fe84f2a
#include <migraphx/gpu/relu.hpp>
#include <migraphx/gpu/context.hpp>
namespace migraphx {
inline namespace MIGRAPHX_INLINE_NS {
namespace gpu {
shape miopen_relu::compute_shape(const std::vector<shape>& inputs) const
{
check_shapes{inputs, *this}.has(2).not_broadcasted();
return inputs.at(1);
}
argument miopen_relu::compute(context& ctx,
const shape& output_shape,
const std::vector<argument>& args) const
{
float alpha = 1;
float beta = 0;
auto x_desc = make_tensor(args[0].get_shape());
auto y_desc = make_tensor(output_shape);
miopenActivationForward(ctx.get_stream().get_miopen(),
ad.get(),
&alpha,
x_desc.get(),
args[0].implicit(),
&beta,
y_desc.get(),
args[1].implicit());
return args[1];
}
} // namespace gpu
} // namespace MIGRAPHX_INLINE_NS
} // namespace migraphx
src/targets/gpu/tanh.cpp deleted 100644 → 0
View file @ 1fe84f2a
#include <migraphx/gpu/tanh.hpp>
#include <migraphx/gpu/context.hpp>
namespace migraphx {
inline namespace MIGRAPHX_INLINE_NS {
namespace gpu {
shape miopen_tanh::compute_shape(const std::vector<shape>& inputs) const
{
check_shapes{inputs, *this}.has(2).packed();
return inputs.at(0);
}
argument miopen_tanh::compute(context& ctx,
const shape& output_shape,
const std::vector<argument>& args) const
{
float alpha = 1;
float beta = 0;
auto x_desc = make_tensor(args[0].get_shape());
auto y_desc = make_tensor(output_shape);
miopenActivationForward(ctx.get_stream().get_miopen(),
ad.get(),
&alpha,
x_desc.get(),
args[0].implicit(),
&beta,
y_desc.get(),
args[1].implicit());
return args[1];
}
} // namespace gpu
} // namespace MIGRAPHX_INLINE_NS
} // namespace migraphx
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment