Merge branch 'master' of https://github.com/ROCmSoftwarePlatform/RTGLib

src/include/rtg/raw_data.hpp

@@ -59,6 +59,7 @@ struct raw_data : raw_data_base
         s.visit_type([&](auto as) { v(make_view(s, as.from(buffer))); });
     }
 
+    /// Returns true if the raw data is only one element
     bool single() const
     {
         auto&& s = static_cast<const Derived&>(*this).get_shape();
@@ -86,17 +87,33 @@ struct raw_data : raw_data_base
         template <class T>
         operator T()
         {
+            assert(self->single());
             return self->template at<T>();
         }
         template <class T>
         operator T*()
         {
-            // TODO: Check type
-            return reinterpret_cast<T*>(self->data());
+            using type = std::remove_cv_t<T>;
+            assert((std::is_void<T>{} or std::is_same<char, type>{} or
+                    std::is_same<unsigned char, type>{} or
+                    self->get_shape().type() == rtg::shape::get_type<T>{}));
+            return reinterpret_cast<type*>(self->data());
         }
     };
 
-    auto_cast get() const { return {static_cast<const Derived*>(this)}; }
+    /// Implicit conversion of raw data pointer
+    auto_cast implicit() const { return {static_cast<const Derived*>(this)}; }
+
+    /// Get a tensor_view to the data
+    template <class T>
+    tensor_view<T> get() const
+    {
+        auto&& s      = static_cast<const Derived&>(*this).get_shape();
+        auto&& buffer = static_cast<const Derived&>(*this).data();
+        if(s.type() != rtg::shape::get_type<T>{})
+            RTG_THROW("Incorrect data type for raw data");
+        return make_view(s, reinterpret_cast<T*>(buffer));
+    }
 };
 
 template <class T,

src/targets/miopen/miopen_target.cpp

@@ -115,31 +115,31 @@ struct miopen_convolution
         float alpha = 1, beta = 0;
         int algo_count;
         miopenConvAlgoPerf_t perf;
-        miopenFindConvolutionForwardAlgorithm(args[0].get(),
+        miopenFindConvolutionForwardAlgorithm(args[0].implicit(),
                                               x_desc.get(),
-                                              args[1].get(),
+                                              args[1].implicit(),
                                               w_desc.get(),
-                                              args[2].get(),
+                                              args[2].implicit(),
                                               cd.get(),
                                               y_desc.get(),
-                                              args[3].get(),
+                                              args[3].implicit(),
                                               1,
                                               &algo_count,
                                               &perf,
                                               nullptr,
                                               0,
                                               false);
-        miopenConvolutionForward(args[0].get(),
+        miopenConvolutionForward(args[0].implicit(),
                                  &alpha,
                                  x_desc.get(),
-                                 args[1].get(),
+                                 args[1].implicit(),
                                  w_desc.get(),
-                                 args[2].get(),
+                                 args[2].implicit(),
                                  cd.get(),
                                  perf.fwd_algo,
                                  &beta,
                                  y_desc.get(),
-                                 args[3].get(),
+                                 args[3].implicit(),
                                  nullptr,
                                  0);
         return args[3];
@@ -161,14 +161,14 @@ struct miopen_relu
         float alpha = 1, beta = 0;
         auto x_desc = make_tensor(args[1].get_shape());
         auto y_desc = make_tensor(output_shape);
-        miopenActivationForward(args[0].get(),
+        miopenActivationForward(args[0].implicit(),
                                 ad.get(),
                                 &alpha,
                                 x_desc.get(),
-                                args[1].get(),
+                                args[1].implicit(),
                                 &beta,
                                 y_desc.get(),
-                                args[2].get());
+                                args[2].implicit());
 
         return args[2];
     }

test/include/verify.hpp

+#ifndef RTG_GUARD_VERIFY_HPP
+#define RTG_GUARD_VERIFY_HPP
+
+#include <algorithm>
+#include <cmath>
+#include <functional>
+#include <iostream>
+#include <numeric>
+
+namespace test {
+
+// Compute the value of a range
+template <class R>
+using range_value = std::decay_t<decltype(*std::declval<R>().begin())>;
+
+struct sum_fn
+{
+    template <class T, class U>
+    auto operator()(T x, U y) const
+    {
+        return x + y;
+    }
+};
+static constexpr sum_fn sum{};
+
+struct max_fn
+{
+    template <class T>
+    static T id(T x)
+    {
+        return x;
+    }
+
+    template <class T, class U>
+    auto operator()(T x, U y) const
+    {
+        return x > y ? x : y;
+    }
+};
+static constexpr max_fn max{};
+
+namespace abs_diff_detail {
+using std::fabs;
+struct fn
+{
+    template <class T, class U>
+    auto operator()(T x, U y) const
+    {
+        return fabs(x - y);
+    }
+};
+
+} // namespace abs_diff_detail
+
+static constexpr abs_diff_detail::fn abs_diff{};
+
+struct not_finite_fn
+{
+    template <class T>
+    bool operator()(T x) const
+    {
+        using std::isfinite;
+        return not isfinite(x);
+    }
+};
+static constexpr not_finite_fn not_finite{};
+
+template <class T, class U>
+T as(T, U x)
+{
+    return x;
+}
+
+struct compare_mag_fn
+{
+    template <class T, class U>
+    bool operator()(T x, U y) const
+    {
+        using std::fabs;
+        return fabs(x) < fabs(y);
+    }
+};
+static constexpr compare_mag_fn compare_mag{};
+
+struct square_diff_fn
+{
+    template <class T, class U>
+    double operator()(T x, U y) const
+    {
+        return (x - y) * (x - y);
+    }
+};
+static constexpr square_diff_fn square_diff{};
+
+template <class R1>
+bool range_empty(R1&& r1)
+{
+    return r1.begin() == r1.end();
+}
+
+template <class R1>
+auto range_distance(R1&& r1)
+{
+    return std::distance(r1.begin(), r1.end());
+}
+
+template <class R1>
+bool range_zero(R1&& r1)
+{
+    return std::all_of(r1.begin(), r1.end(), [](auto x) { return x == 0; });
+}
+
+template <class R1, class R2, class T, class Reducer, class Product>
+T range_product(R1&& r1, R2&& r2, T state, Reducer r, Product p)
+{
+    return std::inner_product(r1.begin(), r1.end(), r2.begin(), state, r, p);
+}
+
+template <class R1, class R2, class Compare>
+std::size_t mismatch_idx(R1&& r1, R2&& r2, Compare compare)
+{
+    auto p = std::mismatch(r1.begin(), r1.end(), r2.begin(), compare);
+    return std::distance(r1.begin(), p.first);
+}
+
+template <class R1, class Predicate>
+long find_idx(R1&& r1, Predicate p)
+{
+    auto it = std::find_if(r1.begin(), r1.end(), p);
+    if(it == r1.end())
+        return -1;
+    else
+        return std::distance(r1.begin(), it);
+}
+
+template <class R1, class R2>
+double max_diff(R1&& r1, R2&& r2)
+{
+    return range_product(r1, r2, 0.0, max, abs_diff);
+}
+
+template <class R1, class R2, class T>
+std::size_t mismatch_diff(R1&& r1, R2&& r2, T diff)
+{
+    return mismatch_idx(r1, r2, [&](auto x, auto y) {
+        auto d = abs_diff(x, y);
+        return !(d > diff && d < diff);
+    });
+}
+
+template <class R1, class R2>
+double rms_range(R1&& r1, R2&& r2)
+{
+    std::size_t n = range_distance(r1);
+    if(n == range_distance(r2))
+    {
+        double square_difference = range_product(r1, r2, 0.0, sum_fn{}, square_diff);
+        double mag1              = *std::max_element(r1.begin(), r1.end(), compare_mag);
+        double mag2              = *std::max_element(r2.begin(), r2.end(), compare_mag);
+        double mag =
+            std::max({std::fabs(mag1), std::fabs(mag2), std::numeric_limits<double>::min()});
+        return std::sqrt(square_difference) / (std::sqrt(n) * mag);
+    }
+    else
+        return std::numeric_limits<range_value<R1>>::max();
+}
+
+template <class R1, class R2>
+bool verify_range(R1&& r1, R2&& r2, double tolerance = 80)
+{
+    double threshold = std::numeric_limits<range_value<R1>>::epsilon() * tolerance;
+    auto error       = rms_range(r1, r2);
+    return error <= threshold;
+}
+} // namespace test
+#endif

test/miopen/miopen.cpp

@@ -10,6 +10,7 @@
 #include <random>
 
 #include "test.hpp"
+#include "verify.hpp"
 
 using hip_ptr       = RTG_MANAGE_PTR(void, hipFree);
 using miopen_handle = RTG_MANAGE_PTR(miopenHandle_t, miopenDestroy);
@@ -92,8 +93,9 @@ std::vector<float> cpu()
     auto x = get_tensor_argument_cpu({rtg::shape::float_type, {4, 3, 3, 3}});
     auto w = get_tensor_argument_cpu({rtg::shape::float_type, {4, 3, 3, 3}});
     p.compile(rtg::cpu::cpu_target{});
-    auto r = p.eval({{"x", x}, {"w", w}});
-    r.visit([&](auto output) { result.assign(output.begin(), output.end()); });
+    auto r      = p.eval({{"x", x}, {"w", w}});
+    auto output = r.get<float>();
+    result.assign(output.begin(), output.end());
     return result;
 }
 
@@ -116,9 +118,7 @@ void test1()
 {
     auto x = cpu();
     auto y = gpu();
-    // TODO: Use expect
-    if(x == y)
-        std::cout << "FAILED" << std::endl;
+    EXPECT(test::verify_range(x, y));
 }
 
 int main() { test1(); }