Merge pull request #73 from ROCmSoftwarePlatform/gpu_op_refactor

Gpu op refactor

src/targets/gpu/CMakeLists.txt

@@ -26,6 +26,14 @@ add_library(migraph_gpu
     hip.cpp
     target.cpp
     lowering.cpp
+    gemm.cpp
+    pooling.cpp
+    convolution.cpp
+    softmax.cpp
+    contiguous.cpp
+    relu.cpp
+    add.cpp
+    batchnorm.cpp
     write_literals.cpp
     rocblas.cpp
 )

src/targets/gpu/add.cpp

+#include <migraph/gpu/add.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+shape hip_add::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_add::compute(context&, const shape&, const std::vector<argument>& args) const
+{
+    device::add(args[2], args[0], args[1]);
+    return args[2];
+}
+
+shape miopen_add::compute_shape(const std::vector<shape>& inputs) const
+{
+    check_shapes{inputs, *this}.has(3).not_broadcasted();
+    return inputs.at(0);
+}
+
+argument miopen_add::compute(context& ctx,
+                             const shape& output_shape,
+                             const std::vector<argument>& args) const
+{
+    float alpha = 1, beta = 0;
+    auto a_desc = make_tensor(args[0].get_shape());
+    auto b_desc = make_tensor(args[1].get_shape());
+    auto c_desc = make_tensor(output_shape);
+    miopenOpTensor(ctx.handle.get(),
+                   miopenTensorOpAdd,
+                   &alpha,
+                   a_desc.get(),
+                   args[0].implicit(),
+                   &alpha,
+                   b_desc.get(),
+                   args[1].implicit(),
+                   &beta,
+                   c_desc.get(),
+                   args[2].implicit());
+    return args[2];
+}
+
+} // namespace gpu
+
+} // namespace migraph

src/targets/gpu/batchnorm.cpp

+#include <migraph/gpu/batchnorm.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+shape miopen_batch_norm_inference::compute_shape(const std::vector<shape>& inputs) const
+{
+    check_shapes{inputs, *this}.has(6);
+    return op.compute_shape({inputs.at(0), inputs.at(1), inputs.at(2), inputs.at(3), inputs.at(4)});
+}
+
+argument miopen_batch_norm_inference::compute(context& ctx,
+                                              const shape& output_shape,
+                                              const std::vector<argument>& args) const
+{
+    auto x_desc  = make_tensor(args[0].get_shape());
+    auto y_desc  = make_tensor(output_shape);
+    auto bn_desc = make_tensor(args[3].get_shape());
+
+    float alpha = 1.0, beta = 0.0f;
+
+    miopenBatchNormalizationForwardInference(ctx.handle.get(),
+                                             miopenBatchNormMode_t(op.bn_mode),
+                                             &alpha,
+                                             &beta,
+                                             x_desc.get(),
+                                             args[0].implicit(),
+                                             y_desc.get(),
+                                             args[5].implicit(),
+                                             bn_desc.get(),
+                                             args[1].implicit(),
+                                             args[2].implicit(),
+                                             args[3].implicit(),
+                                             args[4].implicit(),
+                                             op.epsilon);
+
+    return args[5];
+}
+
+} // namespace gpu
+
+} // namespace migraph

src/targets/gpu/contiguous.cpp

+#include <migraph/gpu/contiguous.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+shape miopen_contiguous::compute_shape(const std::vector<shape>& inputs) const
+{
+    check_shapes{inputs, *this}.has(2);
+    return op.compute_shape({inputs.at(0)});
+}
+argument
+miopen_contiguous::compute(context&, shape output_shape, const std::vector<argument>& args) const
+{
+    assert(output_shape == args[1].get_shape());
+    assert(output_shape.standard());
+    (void)output_shape;
+    device::contiguous(args.at(1), args.at(0));
+    return args.at(1);
+}
+
+} // namespace gpu
+
+} // namespace migraph

src/targets/gpu/convolution.cpp

+#include <migraph/gpu/convolution.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+shape miopen_convolution::compute_shape(const std::vector<shape>& inputs) const
+{
+    check_shapes{inputs, *this}.has(4).standard();
+    return op.compute_shape({inputs.at(0), inputs.at(1)});
+}
+argument miopen_convolution::compute(context& ctx,
+                                     const shape& output_shape,
+                                     const std::vector<argument>& args) const
+{
+    auto x_desc = make_tensor(args[0].get_shape());
+    auto w_desc = make_tensor(args[1].get_shape());
+    auto y_desc = make_tensor(output_shape);
+
+    float alpha = 1, beta = 0;
+    miopenConvolutionForward(ctx.handle.get(),
+                             &alpha,
+                             x_desc.get(),
+                             args[0].implicit(),
+                             w_desc.get(),
+                             args[1].implicit(),
+                             cd.get(),
+                             algo,
+                             &beta,
+                             y_desc.get(),
+                             args[3].implicit(),
+                             args[2].implicit(),
+                             args[2].get_shape().bytes());
+    return args[3];
+}
+
+shape miopen_convolution::compile(context& ctx,
+                                  const shape& output_shape,
+                                  std::vector<instruction_ref> inputs)
+{
+    shape workspace_shape{};
+    auto x_desc = make_tensor(inputs[0]->get_shape());
+    auto w_desc = make_tensor(inputs[1]->get_shape());
+    auto y_desc = make_tensor(output_shape);
+
+    std::size_t workspace_size = 0;
+    miopenConvolutionForwardGetWorkSpaceSize(
+        ctx.handle.get(), w_desc.get(), x_desc.get(), cd.get(), y_desc.get(), &workspace_size);
+    workspace_shape = shape{shape::int8_type, {workspace_size}};
+
+    auto x         = to_gpu(generate_argument(inputs[0]->get_shape()));
+    auto w         = to_gpu(generate_argument(inputs[1]->get_shape()));
+    auto y         = to_gpu(generate_argument(output_shape));
+    auto workspace = allocate_gpu(workspace_shape);
+
+    int algo_count = 1;
+    miopenConvAlgoPerf_t perf;
+    miopenFindConvolutionForwardAlgorithm(ctx.handle.get(),
+                                          x_desc.get(),
+                                          x.implicit(),
+                                          w_desc.get(),
+                                          w.implicit(),
+                                          cd.get(),
+                                          y_desc.get(),
+                                          y.implicit(),
+                                          1,
+                                          &algo_count,
+                                          &perf,
+                                          workspace.implicit(),
+                                          workspace_size,
+                                          false);
+    algo = perf.fwd_algo;
+    return shape{shape::int8_type, {perf.memory}};
+}
+
+} // namespace gpu
+
+} // namespace migraph

src/targets/gpu/gemm.cpp

+#include <migraph/gpu/gemm.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+shape miopen_gemm::compute_shape(const std::vector<shape>& inputs) const
+{
+    check_shapes{inputs, *this}.has(3);
+    return op.compute_shape({inputs.at(0), inputs.at(1)});
+}
+argument miopen_gemm::compute(context& ctx,
+                              const shape& output_shape,
+                              const std::vector<argument>& args) const
+{
+    float alpha     = 1.0f;
+    float beta      = 0.0f;
+    bool transa     = args[0].get_shape().transposed();
+    bool transb     = args[1].get_shape().transposed();
+    rocblas_int lda = args[0].get_shape().strides()[transa ? 1 : 0];
+    rocblas_int ldb = args[1].get_shape().strides()[transb ? 1 : 0];
+    rocblas_int ldc = args[2].get_shape().strides()[0];
+    rocblas_int m   = output_shape.lens()[0];
+    rocblas_int n   = output_shape.lens()[1];
+    rocblas_int k   = args[0].get_shape().lens()[1];
+    rocblas_sgemm(ctx.rbhandle.get(),
+                  transb ? rocblas_operation_transpose : rocblas_operation_none,
+                  transa ? rocblas_operation_transpose : rocblas_operation_none,
+                  n,
+                  m,
+                  k,
+                  &alpha,
+                  args[1].implicit(),
+                  ldb,
+                  args[0].implicit(),
+                  lda,
+                  &beta,
+                  args[2].implicit(),
+                  ldc);
+    return args[2];
+}
+
+} // namespace gpu
+
+} // namespace migraph

src/targets/gpu/include/migraph/gpu/add.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_ADD_HPP
+#define MIGRAPH_GUARD_RTGLIB_ADD_HPP
+
+#include <migraph/gpu/lowering.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/instruction.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/generate.hpp>
+#include <migraph/shape_for_each.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <migraph/gpu/hip.hpp>
+#include <migraph/dfor.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
+#include <migraph/gpu/device/add.hpp>
+#include <migraph/iterator_for.hpp>
+#include <migraph/gpu/rocblas.hpp>
+#include <migraph/gpu/context.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+struct hip_add
+{
+    std::string name() const { return "gpu::add"; }
+    shape compute_shape(const std::vector<shape>& inputs) const;
+    argument compute(context&, const shape&, const std::vector<argument>& args) const;
+};
+
+struct miopen_add
+{
+    std::string name() const { return "gpu::add"; }
+    shape compute_shape(const std::vector<shape>& inputs) const;
+    argument
+    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
+};
+
+} // namespace gpu
+
+} // namespace migraph
+
+#endif

src/targets/gpu/include/migraph/gpu/batchnorm.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_BATCHNORM_HPP
+#define MIGRAPH_GUARD_RTGLIB_BATCHNORM_HPP
+
+#include <migraph/gpu/lowering.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/instruction.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/generate.hpp>
+#include <migraph/shape_for_each.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <migraph/gpu/hip.hpp>
+#include <migraph/dfor.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
+#include <migraph/gpu/device/add.hpp>
+#include <migraph/iterator_for.hpp>
+#include <migraph/gpu/rocblas.hpp>
+#include <migraph/gpu/context.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+struct miopen_batch_norm_inference
+{
+    op::batch_norm_inference op;
+    std::string name() const { return "gpu::batch_norm_inference"; }
+    shape compute_shape(const std::vector<shape>& inputs) const;
+    argument
+    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
+};
+
+} // namespace gpu
+
+} // namespace migraph
+
+#endif

src/targets/gpu/include/migraph/gpu/contiguous.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_CONTIGUOUS_HPP
+#define MIGRAPH_GUARD_RTGLIB_CONTIGUOUS_HPP
+
+#include <migraph/gpu/lowering.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/instruction.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/generate.hpp>
+#include <migraph/shape_for_each.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <migraph/gpu/hip.hpp>
+#include <migraph/dfor.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
+#include <migraph/gpu/device/add.hpp>
+#include <migraph/iterator_for.hpp>
+#include <migraph/gpu/rocblas.hpp>
+#include <migraph/gpu/context.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+struct miopen_contiguous
+{
+    op::contiguous op;
+    std::string name() const { return "gpu::contiguous"; }
+    shape compute_shape(const std::vector<shape>& inputs) const;
+    argument compute(context&, shape output_shape, const std::vector<argument>& args) const;
+};
+
+} // namespace gpu
+
+} // namespace migraph
+
+#endif

src/targets/gpu/include/migraph/gpu/convolution.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_CONVOLUTION_HPP
+#define MIGRAPH_GUARD_RTGLIB_CONVOLUTION_HPP
+
+#include <migraph/gpu/lowering.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/instruction.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/generate.hpp>
+#include <migraph/shape_for_each.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <migraph/gpu/hip.hpp>
+#include <migraph/dfor.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
+#include <migraph/gpu/device/add.hpp>
+#include <migraph/iterator_for.hpp>
+#include <migraph/gpu/rocblas.hpp>
+#include <migraph/gpu/context.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+struct miopen_convolution
+{
+    op::convolution op;
+    shared<convolution_descriptor> cd;
+    miopenConvFwdAlgorithm_t algo{};
+
+    std::string name() const { return "gpu::convolution"; }
+    shape compute_shape(const std::vector<shape>& inputs) const;
+    argument
+    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
+    shape compile(context& ctx, const shape& output_shape, std::vector<instruction_ref> inputs);
+    friend std::ostream& operator<<(std::ostream& os, const miopen_convolution& self)
+    {
+        os << self.name() << "[";
+        os << self.op << ", ";
+        os << "algo=" << self.algo;
+        os << "]";
+        return os;
+    }
+};
+
+} // namespace gpu
+
+} // namespace migraph
+
+#endif

src/targets/gpu/include/migraph/gpu/gemm.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_GEMM_HPP
+#define MIGRAPH_GUARD_RTGLIB_GEMM_HPP
+
+#include <migraph/gpu/lowering.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/instruction.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/generate.hpp>
+#include <migraph/shape_for_each.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <migraph/gpu/hip.hpp>
+#include <migraph/dfor.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
+#include <migraph/gpu/device/add.hpp>
+#include <migraph/iterator_for.hpp>
+#include <migraph/gpu/rocblas.hpp>
+#include <migraph/gpu/context.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+struct miopen_gemm
+{
+    op::gemm op;
+    std::string name() const { return "gpu::gemm"; }
+    shape compute_shape(const std::vector<shape>& inputs) const;
+    argument
+    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
+};
+
+} // namespace gpu
+
+} // namespace migraph
+
+#endif

src/targets/gpu/include/migraph/gpu/pooling.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_POOLING_HPP
+#define MIGRAPH_GUARD_RTGLIB_POOLING_HPP
+
+#include <migraph/gpu/lowering.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/instruction.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/generate.hpp>
+#include <migraph/shape_for_each.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <migraph/gpu/hip.hpp>
+#include <migraph/dfor.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
+#include <migraph/gpu/device/add.hpp>
+#include <migraph/iterator_for.hpp>
+#include <migraph/gpu/rocblas.hpp>
+#include <migraph/gpu/context.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+struct miopen_pooling
+{
+    op::pooling op;
+    shared<pooling_descriptor> pd;
+
+    std::string name() const { return "gpu::pooling"; }
+    shape compute_shape(const std::vector<shape>& inputs) const;
+    argument
+    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
+};
+
+} // namespace gpu
+
+} // namespace migraph
+
+#endif

src/targets/gpu/include/migraph/gpu/relu.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_RELU_HPP
+#define MIGRAPH_GUARD_RTGLIB_RELU_HPP
+
+#include <migraph/gpu/lowering.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/instruction.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/generate.hpp>
+#include <migraph/shape_for_each.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <migraph/gpu/hip.hpp>
+#include <migraph/dfor.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
+#include <migraph/gpu/device/add.hpp>
+#include <migraph/iterator_for.hpp>
+#include <migraph/gpu/rocblas.hpp>
+#include <migraph/gpu/context.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+struct miopen_relu
+{
+    shared<activation_descriptor> ad;
+    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;
+};
+
+} // namespace gpu
+
+} // namespace migraph
+
+#endif

src/targets/gpu/include/migraph/gpu/softmax.hpp

+#ifndef MIGRAPH_GUARD_RTGLIB_SOFTMAX_HPP
+#define MIGRAPH_GUARD_RTGLIB_SOFTMAX_HPP
+
+#include <migraph/gpu/lowering.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/instruction.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/generate.hpp>
+#include <migraph/shape_for_each.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <migraph/gpu/hip.hpp>
+#include <migraph/dfor.hpp>
+#include <migraph/gpu/device/contiguous.hpp>
+#include <migraph/gpu/device/add.hpp>
+#include <migraph/iterator_for.hpp>
+#include <migraph/gpu/rocblas.hpp>
+#include <migraph/gpu/context.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+struct miopen_softmax
+{
+    op::softmax op;
+    std::string name() const { return "gpu::softmax"; }
+    shape compute_shape(const std::vector<shape>& inputs) const;
+    argument
+    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const;
+};
+
+} // namespace gpu
+
+} // namespace migraph
+
+#endif

src/targets/gpu/lowering.cpp

@@ -13,337 +13,19 @@
 #include <migraph/iterator_for.hpp>
 #include <migraph/gpu/rocblas.hpp>
 #include <migraph/gpu/context.hpp>
+#include <migraph/gpu/convolution.hpp>
+#include <migraph/gpu/contiguous.hpp>
+#include <migraph/gpu/relu.hpp>
+#include <migraph/gpu/softmax.hpp>
+#include <migraph/gpu/add.hpp>
+#include <migraph/gpu/batchnorm.hpp>
+#include <migraph/gpu/pooling.hpp>
+#include <migraph/gpu/gemm.hpp>
 #include <utility>
 
 namespace migraph {
 namespace gpu {
 
-struct miopen_batch_norm_inference
-{
-    op::batch_norm_inference op;
-
-    std::string name() const { return "gpu::batch_norm_inference"; }
-
-    shape compute_shape(const std::vector<shape>& inputs) const
-    {
-        check_shapes{inputs, *this}.has(6);
-        return op.compute_shape(
-            {inputs.at(0), inputs.at(1), inputs.at(2), inputs.at(3), inputs.at(4)});
-    }
-
-    argument
-    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const
-    {
-        auto x_desc  = make_tensor(args[0].get_shape());
-        auto y_desc  = make_tensor(output_shape);
-        auto bn_desc = make_tensor(args[3].get_shape());
-
-        float alpha = 1.0, beta = 0.0f;
-
-        miopenBatchNormalizationForwardInference(ctx.handle.get(),
-                                                 miopenBatchNormMode_t(op.bn_mode),
-                                                 &alpha,
-                                                 &beta,
-                                                 x_desc.get(),
-                                                 args[0].implicit(),
-                                                 y_desc.get(),
-                                                 args[5].implicit(),
-                                                 bn_desc.get(),
-                                                 args[1].implicit(),
-                                                 args[2].implicit(),
-                                                 args[3].implicit(),
-                                                 args[4].implicit(),
-                                                 op.epsilon);
-
-        return args[5];
-    }
-};
-
-struct miopen_convolution
-{
-    op::convolution op;
-    shared<convolution_descriptor> cd;
-    miopenConvFwdAlgorithm_t algo{};
-
-    std::string name() const { return "gpu::convolution"; }
-    shape compute_shape(const std::vector<shape>& inputs) const
-    {
-        check_shapes{inputs, *this}.has(4).standard();
-        return op.compute_shape({inputs.at(0), inputs.at(1)});
-    }
-    argument
-    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const
-    {
-        auto x_desc = make_tensor(args[0].get_shape());
-        auto w_desc = make_tensor(args[1].get_shape());
-        auto y_desc = make_tensor(output_shape);
-
-        float alpha = 1, beta = 0;
-        miopenConvolutionForward(ctx.handle.get(),
-                                 &alpha,
-                                 x_desc.get(),
-                                 args[0].implicit(),
-                                 w_desc.get(),
-                                 args[1].implicit(),
-                                 cd.get(),
-                                 algo,
-                                 &beta,
-                                 y_desc.get(),
-                                 args[3].implicit(),
-                                 args[2].implicit(),
-                                 args[2].get_shape().bytes());
-        return args[3];
-    }
-
-    shape compile(context& ctx, const shape& output_shape, std::vector<instruction_ref> inputs)
-    {
-        shape workspace_shape{};
-        auto x_desc = make_tensor(inputs[0]->get_shape());
-        auto w_desc = make_tensor(inputs[1]->get_shape());
-        auto y_desc = make_tensor(output_shape);
-
-        std::size_t workspace_size = 0;
-        miopenConvolutionForwardGetWorkSpaceSize(
-            ctx.handle.get(), w_desc.get(), x_desc.get(), cd.get(), y_desc.get(), &workspace_size);
-        workspace_shape = shape{shape::int8_type, {workspace_size}};
-
-        auto x = to_gpu(generate_argument(inputs[0]->get_shape()));
-        gpu_sync();
-        auto w = to_gpu(generate_argument(inputs[1]->get_shape()));
-        gpu_sync();
-        auto y = to_gpu(generate_argument(output_shape));
-        gpu_sync();
-        auto workspace = allocate_gpu(workspace_shape);
-
-        int algo_count = 1;
-        miopenConvAlgoPerf_t perf;
-        miopenFindConvolutionForwardAlgorithm(ctx.handle.get(),
-                                              x_desc.get(),
-                                              x.implicit(),
-                                              w_desc.get(),
-                                              w.implicit(),
-                                              cd.get(),
-                                              y_desc.get(),
-                                              y.implicit(),
-                                              1,
-                                              &algo_count,
-                                              &perf,
-                                              workspace.implicit(),
-                                              workspace_size,
-                                              false);
-        algo = perf.fwd_algo;
-        return shape{shape::int8_type, {perf.memory}};
-    }
-
-    friend std::ostream& operator<<(std::ostream& os, const miopen_convolution& self)
-    {
-        os << self.name() << "[";
-        os << self.op << ", ";
-        os << "algo=" << self.algo;
-        os << "]";
-        return os;
-    }
-};
-
-struct miopen_pooling
-{
-    op::pooling op;
-    shared<pooling_descriptor> pd;
-
-    std::string name() const { return "gpu::pooling"; }
-    shape compute_shape(const std::vector<shape>& inputs) const
-    {
-        check_shapes{inputs, *this}.has(2).standard();
-        return op.compute_shape({inputs.at(0)});
-    }
-    argument
-    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const
-    {
-        auto x_desc = make_tensor(args[0].get_shape());
-        auto y_desc = make_tensor(output_shape);
-
-        float alpha = 1, beta = 0;
-
-        miopenPoolingForward(ctx.handle.get(),
-                             pd.get(),
-                             &alpha,
-                             x_desc.get(),
-                             args[0].implicit(),
-                             &beta,
-                             y_desc.get(),
-                             args[1].implicit(),
-                             false,
-                             nullptr,
-                             0);
-
-        return args[1];
-    }
-};
-
-struct hip_add
-{
-    std::string name() const { return "gpu::add"; }
-    shape 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 compute(context&, const shape&, const std::vector<argument>& args) const
-    {
-        device::add(args[2], args[0], args[1]);
-        return args[2];
-    }
-};
-
-struct miopen_add
-{
-    std::string name() const { return "gpu::add"; }
-    shape compute_shape(const std::vector<shape>& inputs) const
-    {
-        check_shapes{inputs, *this}.has(3).not_broadcasted();
-        return inputs.at(0);
-    }
-
-    argument
-    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const
-    {
-        float alpha = 1, beta = 0;
-        auto a_desc = make_tensor(args[0].get_shape());
-        auto b_desc = make_tensor(args[1].get_shape());
-        auto c_desc = make_tensor(output_shape);
-        miopenOpTensor(ctx.handle.get(),
-                       miopenTensorOpAdd,
-                       &alpha,
-                       a_desc.get(),
-                       args[0].implicit(),
-                       &alpha,
-                       b_desc.get(),
-                       args[1].implicit(),
-                       &beta,
-                       c_desc.get(),
-                       args[2].implicit());
-        return args[2];
-    }
-};
-
-struct miopen_gemm
-{
-    op::gemm op;
-    std::string name() const { return "gpu::gemm"; }
-    shape compute_shape(const std::vector<shape>& inputs) const
-    {
-        check_shapes{inputs, *this}.has(3);
-        return op.compute_shape({inputs.at(0), inputs.at(1)});
-    }
-    argument
-    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const
-    {
-        float alpha     = 1.0f;
-        float beta      = 0.0f;
-        bool transa     = args[0].get_shape().transposed();
-        bool transb     = args[1].get_shape().transposed();
-        rocblas_int lda = args[0].get_shape().strides()[transa ? 1 : 0];
-        rocblas_int ldb = args[1].get_shape().strides()[transb ? 1 : 0];
-        rocblas_int ldc = args[2].get_shape().strides()[0];
-        rocblas_int m   = output_shape.lens()[0];
-        rocblas_int n   = output_shape.lens()[1];
-        rocblas_int k   = args[0].get_shape().lens()[1];
-        rocblas_sgemm(ctx.rbhandle.get(),
-                      transb ? rocblas_operation_transpose : rocblas_operation_none,
-                      transa ? rocblas_operation_transpose : rocblas_operation_none,
-                      n,
-                      m,
-                      k,
-                      &alpha,
-                      args[1].implicit(),
-                      ldb,
-                      args[0].implicit(),
-                      lda,
-                      &beta,
-                      args[2].implicit(),
-                      ldc);
-        return args[2];
-    }
-};
-
-struct miopen_contiguous
-{
-    op::contiguous op;
-    std::string name() const { return "gpu::contiguous"; }
-    shape compute_shape(const std::vector<shape>& inputs) const
-    {
-        check_shapes{inputs, *this}.has(2);
-        return op.compute_shape({inputs.at(0)});
-    }
-    argument compute(context&, shape output_shape, const std::vector<argument>& args) const
-    {
-        assert(output_shape == args[1].get_shape());
-        assert(output_shape.standard());
-        (void)output_shape;
-        device::contiguous(args.at(1), args.at(0));
-        return args.at(1);
-    }
-};
-
-struct miopen_relu
-{
-    shared<activation_descriptor> ad;
-    std::string name() const { return "gpu::relu"; }
-    shape compute_shape(const std::vector<shape>& inputs) const
-    {
-        check_shapes{inputs, *this}.has(2).not_broadcasted();
-        return inputs.at(1);
-    }
-
-    argument
-    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const
-    {
-        float alpha = 1, beta = 0;
-        auto x_desc = make_tensor(args[0].get_shape());
-        auto y_desc = make_tensor(output_shape);
-        miopenActivationForward(ctx.handle.get(),
-                                ad.get(),
-                                &alpha,
-                                x_desc.get(),
-                                args[0].implicit(),
-                                &beta,
-                                y_desc.get(),
-                                args[1].implicit());
-
-        return args[1];
-    }
-};
-
-struct miopen_softmax
-{
-    op::softmax op;
-    std::string name() const { return "gpu::softmax"; }
-    shape compute_shape(const std::vector<shape>& inputs) const
-    {
-        check_shapes{inputs, *this}.has(2).standard();
-        return op.compute_shape({inputs.at(0)});
-    }
-
-    argument
-    compute(context& ctx, const shape& output_shape, const std::vector<argument>& args) const
-    {
-        float alpha = 1, beta = 0;
-        auto x_desc = make_tensor(args[0].get_shape());
-        auto y_desc = make_tensor(output_shape);
-        miopenSoftmaxForward(ctx.handle.get(),
-                             &alpha,
-                             x_desc.get(),
-                             args[0].implicit(),
-                             &beta,
-                             y_desc.get(),
-                             args[1].implicit());
-
-        return args[1];
-    }
-};
-
 struct miopen_apply
 {
     program* prog = nullptr;

src/targets/gpu/pooling.cpp

+#include <migraph/gpu/pooling.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+shape miopen_pooling::compute_shape(const std::vector<shape>& inputs) const
+{
+    check_shapes{inputs, *this}.has(2).standard();
+    return op.compute_shape({inputs.at(0)});
+}
+argument miopen_pooling::compute(context& ctx,
+                                 const shape& output_shape,
+                                 const std::vector<argument>& args) const
+{
+    auto x_desc = make_tensor(args[0].get_shape());
+    auto y_desc = make_tensor(output_shape);
+
+    float alpha = 1, beta = 0;
+
+    miopenPoolingForward(ctx.handle.get(),
+                         pd.get(),
+                         &alpha,
+                         x_desc.get(),
+                         args[0].implicit(),
+                         &beta,
+                         y_desc.get(),
+                         args[1].implicit(),
+                         false,
+                         nullptr,
+                         0);
+
+    return args[1];
+}
+
+} // namespace gpu
+
+} // namespace migraph

src/targets/gpu/relu.cpp

+#include <migraph/gpu/relu.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <utility>
+
+namespace migraph {
+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, beta = 0;
+    auto x_desc = make_tensor(args[0].get_shape());
+    auto y_desc = make_tensor(output_shape);
+    miopenActivationForward(ctx.handle.get(),
+                            ad.get(),
+                            &alpha,
+                            x_desc.get(),
+                            args[0].implicit(),
+                            &beta,
+                            y_desc.get(),
+                            args[1].implicit());
+
+    return args[1];
+}
+
+} // namespace gpu
+
+} // namespace migraph

src/targets/gpu/softmax.cpp

+#include <migraph/gpu/softmax.hpp>
+#include <migraph/operators.hpp>
+#include <migraph/manage_ptr.hpp>
+#include <migraph/gpu/miopen.hpp>
+#include <utility>
+
+namespace migraph {
+namespace gpu {
+
+shape miopen_softmax::compute_shape(const std::vector<shape>& inputs) const
+{
+    check_shapes{inputs, *this}.has(2).standard();
+    return op.compute_shape({inputs.at(0)});
+}
+
+argument miopen_softmax::compute(context& ctx,
+                                 const shape& output_shape,
+                                 const std::vector<argument>& args) const
+{
+    float alpha = 1, beta = 0;
+    auto x_desc = make_tensor(args[0].get_shape());
+    auto y_desc = make_tensor(output_shape);
+    miopenSoftmaxForward(ctx.handle.get(),
+                         &alpha,
+                         x_desc.get(),
+                         args[0].implicit(),
+                         &beta,
+                         y_desc.get(),
+                         args[1].implicit());
+
+    return args[1];
+}
+
+} // namespace gpu
+
+} // namespace migraph