Refactored the rest of the GPU operators into individual translation units

src/targets/gpu/CMakeLists.txt

@@ -29,6 +29,11 @@ add_library(migraph_gpu
     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/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
+{
+    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
+{
+    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/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
+{
+    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

@@ -14,6 +14,11 @@
 #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>
@@ -21,171 +26,6 @@
 namespace migraph {
 namespace gpu {
 
-struct miopen_batch_norm_inference
-{
-    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 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_contiguous
-{
-    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
-{
-    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/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