Refactor hip functions to header files

src/targets/miopen/include/rtg/miopen/hip.hpp

+#ifndef RTG_GUARD_RTGLIB_HIP_HPP
+#define RTG_GUARD_RTGLIB_HIP_HPP
+
+#include <rtg/manage_ptr.hpp>
+
+#include <miopen/miopen.h>
+
+namespace rtg { namespace miopen {
+
+using hip_ptr       = RTG_MANAGE_PTR(void, hipFree);
+
+inline hip_ptr gpu_allocate(std::size_t sz)
+{
+    void* result;
+    // TODO: Check status
+    hipMalloc(&result, sz);
+    return hip_ptr{result};
+}
+
+template <class T>
+hip_ptr write_to_gpu(const T& x)
+{
+    using type  = typename T::value_type;
+    auto size   = x.size() * sizeof(type);
+    auto result = gpu_allocate(size);
+    // TODO: Check status
+    hipMemcpy(result.get(), x.data(), size, hipMemcpyHostToDevice);
+    return result;
+}
+
+template <class T>
+std::vector<T> read_from_gpu(const void* x, std::size_t sz)
+{
+    std::vector<T> result(sz);
+    // TODO: Check status
+    hipMemcpy(result.data(), x, sz * sizeof(T), hipMemcpyDeviceToHost);
+    return result;
+}
+
+struct hip_allocate
+{
+    std::string name() const { return "hip::allocate"; }
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        check_shapes{inputs}.has(1);
+        return inputs.front();
+    }
+    argument compute(shape output_shape, std::vector<argument>) const
+    {
+        char* data = nullptr;
+        // TODO: Check return status
+        hipMalloc(&data, output_shape.bytes());
+        return {output_shape, data};
+    }
+};
+
+struct hip_free
+{
+    std::string name() const { return "hip::free"; }
+    shape compute_shape(std::vector<shape> inputs) const
+    {
+        check_shapes{inputs}.has(1);
+        return {};
+    }
+    argument compute(shape, std::vector<argument> args) const
+    {
+        // TODO: Check return status
+        hipFree(args.front().data());
+        return {};
+    }
+};
+
+} // namespace miopen
+
+} // namespace rtg
+
+#endif

src/targets/miopen/include/rtg/miopen/miopen.hpp

+#ifndef RTG_GUARD_RTGLIB_MIOPEN_HPP
+#define RTG_GUARD_RTGLIB_MIOPEN_HPP
+
+#include <rtg/manage_ptr.hpp>
+
+#include <miopen/miopen.h>
+
+namespace rtg { namespace miopen {
+
+using miopen_handle     = RTG_MANAGE_PTR(miopenHandle_t, miopenDestroy);
+using tensor_descriptor = RTG_MANAGE_PTR(miopenTensorDescriptor_t, miopenDestroyTensorDescriptor);
+using convolution_descriptor = RTG_MANAGE_PTR(miopenConvolutionDescriptor_t,
+                                              miopenDestroyConvolutionDescriptor);
+using activation_descriptor  = RTG_MANAGE_PTR(miopenActivationDescriptor_t,
+                                             miopenDestroyActivationDescriptor);
+
+template <class Result, class F, class... Ts>
+Result make_obj(F f, Ts... xs)
+{
+    typename Result::pointer x = nullptr;
+    auto status                = f(&x, xs...);
+    Result r{x};
+    if(status != miopenStatusSuccess)
+        RTG_THROW("MIOpen call failed");
+    return r;
+}
+
+inline tensor_descriptor make_tensor(const rtg::shape& s)
+{
+    auto t = make_obj<tensor_descriptor>(&miopenCreateTensorDescriptor);
+    // Convert to ints
+    std::vector<int> lens(s.lens().begin(), s.lens().end());
+    std::vector<int> strides(s.strides().begin(), s.strides().end());
+    miopenDataType_t d;
+    if(s.type() == shape::float_type)
+        d = miopenFloat;
+    else
+        RTG_THROW("Unsupported type");
+    miopenSetTensorDescriptor(t.get(), d, s.lens().size(), lens.data(), strides.data());
+    return t;
+}
+
+inline convolution_descriptor make_conv(const rtg::convolution& op)
+{
+    auto c = make_obj<convolution_descriptor>(&miopenCreateConvolutionDescriptor);
+    miopenInitConvolutionDescriptor(c.get(),
+                                    miopenConvolution,
+                                    op.padding[0],
+                                    op.padding[1],
+                                    op.stride[0],
+                                    op.stride[1],
+                                    op.dilation[0],
+                                    op.dilation[1]);
+    return c;
+}
+
+inline activation_descriptor make_relu()
+{
+    auto ad = make_obj<activation_descriptor>(&miopenCreateActivationDescriptor);
+    miopenSetActivationDescriptor(ad.get(), miopenActivationRELU, 0, 0, 0);
+    return ad;
+}
+
+} // namespace miopen
+
+} // namespace rtg
+
+#endif

src/targets/miopen/miopen_target.cpp

@@ -2,99 +2,12 @@
 #include <rtg/manage_ptr.hpp>
 #include <rtg/instruction.hpp>
 #include <rtg/operators.hpp>
-
-#include <miopen/miopen.h>
+#include <rtg/miopen/miopen.hpp>
+#include <rtg/miopen/hip.hpp>
 
 namespace rtg {
 namespace miopen {
 
-struct hip_allocate
-{
-    std::string name() const { return "hip::allocate"; }
-    shape compute_shape(std::vector<shape> inputs) const
-    {
-        check_shapes{inputs}.has(1);
-        return inputs.front();
-    }
-    argument compute(shape output_shape, std::vector<argument>) const
-    {
-        char* data = nullptr;
-        // TODO: Check return status
-        hipMalloc(&data, output_shape.bytes());
-        return {output_shape, data};
-    }
-};
-
-struct hip_free
-{
-    std::string name() const { return "hip::free"; }
-    shape compute_shape(std::vector<shape> inputs) const
-    {
-        check_shapes{inputs}.has(1);
-        return {};
-    }
-    argument compute(shape, std::vector<argument> args) const
-    {
-        // TODO: Check return status
-        hipFree(args.front().data());
-        return {};
-    }
-};
-
-using miopen_handle     = RTG_MANAGE_PTR(miopenHandle_t, miopenDestroy);
-using tensor_descriptor = RTG_MANAGE_PTR(miopenTensorDescriptor_t, miopenDestroyTensorDescriptor);
-using convolution_descriptor = RTG_MANAGE_PTR(miopenConvolutionDescriptor_t,
-                                              miopenDestroyConvolutionDescriptor);
-using activation_descriptor  = RTG_MANAGE_PTR(miopenActivationDescriptor_t,
-                                             miopenDestroyActivationDescriptor);
-
-template <class Result, class F, class... Ts>
-Result make_obj(F f, Ts... xs)
-{
-    typename Result::pointer x = nullptr;
-    auto status                = f(&x, xs...);
-    Result r{x};
-    if(status != miopenStatusSuccess)
-        RTG_THROW("MIOpen call failed");
-    return r;
-}
-
-tensor_descriptor make_tensor(const rtg::shape& s)
-{
-    auto t = make_obj<tensor_descriptor>(&miopenCreateTensorDescriptor);
-    // Convert to ints
-    std::vector<int> lens(s.lens().begin(), s.lens().end());
-    std::vector<int> strides(s.strides().begin(), s.strides().end());
-    miopenDataType_t d;
-    if(s.type() == shape::float_type)
-        d = miopenFloat;
-    else
-        RTG_THROW("Unsupported type");
-    miopenSetTensorDescriptor(t.get(), d, s.lens().size(), lens.data(), strides.data());
-    return t;
-}
-
-convolution_descriptor make_conv(const rtg::convolution& op)
-{
-    auto c = make_obj<convolution_descriptor>(&miopenCreateConvolutionDescriptor);
-    miopenInitConvolutionDescriptor(c.get(),
-                                    miopenConvolution,
-                                    op.padding[0],
-                                    op.padding[1],
-                                    op.stride[0],
-                                    op.stride[1],
-                                    op.dilation[0],
-                                    op.dilation[1]);
-    return c;
-}
-
-activation_descriptor make_relu()
-{
-    auto ad = make_obj<activation_descriptor>(&miopenCreateActivationDescriptor);
-    miopenSetActivationDescriptor(ad.get(), miopenActivationRELU, 0, 0, 0);
-    return ad;
-}
-
 struct miopen_convolution
 {
     convolution op;

test/miopen/miopen.cpp

@@ -4,6 +4,8 @@
 #include <rtg/generate.hpp>
 #include <rtg/cpu/cpu_target.hpp>
 #include <rtg/miopen/miopen_target.hpp>
+#include <rtg/miopen/miopen.hpp>
+#include <rtg/miopen/hip.hpp>
 #include <rtg/manage_ptr.hpp>
 
 #include <miopen/miopen.h>
@@ -11,48 +13,6 @@
 #include "test.hpp"
 #include "verify.hpp"
 
-using hip_ptr       = RTG_MANAGE_PTR(void, hipFree);
-using miopen_handle = RTG_MANAGE_PTR(miopenHandle_t, miopenDestroy);
-
-template <class Result, class F, class... Ts>
-Result make_obj(F f, Ts... xs)
-{
-    typename Result::pointer x = nullptr;
-    auto status                = f(&x, xs...);
-    Result r{x};
-    if(status != miopenStatusSuccess)
-        RTG_THROW("MIOpen call failed");
-    return r;
-}
-
-hip_ptr hip_allocate(std::size_t sz)
-{
-    void* result;
-    // TODO: Check status
-    hipMalloc(&result, sz);
-    return hip_ptr{result};
-}
-
-template <class T>
-hip_ptr write(const T& x)
-{
-    using type  = typename T::value_type;
-    auto size   = x.size() * sizeof(type);
-    auto result = hip_allocate(size);
-    // TODO: Check status
-    hipMemcpy(result.get(), x.data(), size, hipMemcpyHostToDevice);
-    return result;
-}
-
-template <class T>
-std::vector<T> read(const void* x, std::size_t sz)
-{
-    std::vector<T> result(sz);
-    // TODO: Check status
-    hipMemcpy(result.data(), x, sz * sizeof(T), hipMemcpyDeviceToHost);
-    return result;
-}
-
 rtg::program create_program()
 {
     rtg::program p;
@@ -67,7 +27,7 @@ rtg::program create_program()
 rtg::argument get_tensor_argument_gpu(rtg::shape s)
 {
     auto v = rtg::generate_tensor_data<float>(s);
-    auto p = rtg::share(write(v));
+    auto p = rtg::share(rtg::miopen::write_to_gpu(v));
     return {s, [p]() mutable { return reinterpret_cast<char*>(p.get()); }};
 }
 
@@ -92,10 +52,10 @@ std::vector<float> gpu()
     auto w = get_tensor_argument_gpu({rtg::shape::float_type, {4, 3, 3, 3}});
     p.compile(rtg::miopen::miopen_target{});
     auto y      = get_tensor_argument_gpu(p.get_parameter_shape("output"));
-    auto handle = make_obj<miopen_handle>(&miopenCreate);
+    auto handle = rtg::miopen::make_obj<rtg::miopen::miopen_handle>(&miopenCreate);
     auto r      = p.eval(
         {{"x", x}, {"w", w}, {"output", y}, {"handle", {rtg::shape::any_type, handle.get()}}});
-    result = read<float>(r.data(), r.get_shape().elements());
+    result = rtg::miopen::read_from_gpu<float>(r.data(), r.get_shape().elements());
     return result;
 }