Initial commit

src/targets/gpu/CMakeLists.txt

@@ -26,6 +26,9 @@ add_library(migraph_gpu
     hip.cpp
     target.cpp
     lowering.cpp
+    gemm.cpp
+    pooling.cpp
+    convolution.cpp
     write_literals.cpp
     rocblas.cpp
 )

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/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
+{
+    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
+{
+    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
+{
+    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/lowering.cpp

@@ -13,6 +13,9 @@
 #include <migraph/iterator_for.hpp>
 #include <migraph/gpu/rocblas.hpp>
 #include <migraph/gpu/context.hpp>
+#include <migraph/gpu/convolution.hpp>
+#include <migraph/gpu/pooling.hpp>
+#include <migraph/gpu/gemm.hpp>
 #include <utility>
 
 namespace migraph {
@@ -59,124 +62,6 @@ struct miopen_batch_norm_inference
     }
 };
 
-struct miopen_convolution
-{
-    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()));
-        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}};
-    }
-
-    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
-{
-    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"; }
@@ -225,46 +110,6 @@ struct miopen_add
     }
 };
 
-struct miopen_gemm
-{
-    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
 {
     contiguous op;

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
+