Merge branch 'develop' into register_target

ROCm · May 27, 2024 · 5e3411f · 5e3411f
2 parents 9d8bcf8 + 2dcca47
commit 5e3411f
Show file tree

Hide file tree

Showing 4 changed files with 19 additions and 6 deletions.
diff --git a/requirements.txt b/requirements.txt
@@ -24,7 +24,7 @@
 google/protobuf@v3.19.0 -DCMAKE_POSITION_INDEPENDENT_CODE=On -X subdir -Dprotobuf_BUILD_TESTS=Off
 nlohmann/json@v3.8.0
 ROCm/half@rocm-5.6.0
-pybind/pybind11@d159a563383d10c821ba7b2a71905d1207db6de4 --build
+pybind/pybind11@3e9dfa2866941655c56877882565e7577de6fc7b --build
 msgpack/msgpack-c@cpp-3.3.0 -DMSGPACK_BUILD_TESTS=Off
 sqlite3@3.43.2 -DCMAKE_POSITION_INDEPENDENT_CODE=On
 ROCm/composable_kernel@57cdd70b7cb14e5e3b60cd9a5f96ba8dc343763e -DCK_BUILD_JIT_LIB=On -DCMAKE_POSITION_INDEPENDENT_CODE=On

diff --git a/src/include/migraphx/op/mod.hpp b/src/include/migraphx/op/mod.hpp
@@ -1,7 +1,7 @@
 /*
  * The MIT License (MIT)
  *
- * Copyright (c) 2015-2022 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -38,7 +38,6 @@ struct mod : binary<mod>
     {
         auto a           = base_attributes();
         a["commutative"] = false;
-        a["point_op"]    = "${function:fmod}((${function:remainder}(${0}, ${1})) + ${1}, ${1})";
         return a;
     }
     auto apply() const

diff --git a/src/targets/gpu/kernels/include/migraphx/kernels/math.hpp b/src/targets/gpu/kernels/include/migraphx/kernels/math.hpp
@@ -1,7 +1,7 @@
 /*
  * The MIT License (MIT)
  *
- * Copyright (c) 2015-2023 Advanced Micro Devices, Inc. All rights reserved.
+ * Copyright (c) 2015-2024 Advanced Micro Devices, Inc. All rights reserved.
  *
  * Permission is hereby granted, free of charge, to any person obtaining a copy
  * of this software and associated documentation files (the "Software"), to deal
@@ -256,6 +256,21 @@ constexpr auto min(const T& a, const U& b)
     return min<common_type_t<T, U>>(a, b);
 }
 
+template <class T, MIGRAPHX_REQUIRES(not is_any_vec<T>())>
+constexpr T mod(const T& a, const T& b)
+{
+    if constexpr(is_integral<T>{})
+        // onnx mod operator requires numpy style modulus
+        return ((a % b) + b) % b;
+    return static_cast<T>(fmod(remainder(a, b) + b, b));
+}
+
+template <class T, class U, MIGRAPHX_REQUIRES(not is_same<T, U>{} and not is_any_vec<T, U>())>
+constexpr auto mod(const T& a, const U& b)
+{
+    return mod<common_type_t<T, U>>(a, b);
+}
+
 MIGRAPHX_DEVICE_MATH_VEC(abs)
 MIGRAPHX_DEVICE_MATH_VEC(acos)
 MIGRAPHX_DEVICE_MATH_VEC(acosh)
@@ -275,6 +290,7 @@ MIGRAPHX_DEVICE_MATH_VEC(isnan)
 MIGRAPHX_DEVICE_MATH_VEC(log)
 MIGRAPHX_DEVICE_MATH_VEC(max)
 MIGRAPHX_DEVICE_MATH_VEC(min)
+MIGRAPHX_DEVICE_MATH_VEC(mod)
 MIGRAPHX_DEVICE_MATH_VEC(nearbyint)
 MIGRAPHX_DEVICE_MATH_VEC(pow)
 MIGRAPHX_DEVICE_MATH_VEC(remainder)

diff --git a/test/py/onnx_backend_test.py b/test/py/onnx_backend_test.py
@@ -119,8 +119,6 @@ def disabled_tests_onnx_1_7_0(backend_test):
     backend_test.exclude(r'test_einsum_transpose_cpu')
     backend_test.exclude(r'test_maxunpool_export_with_output_shape_cpu')
     backend_test.exclude(r'test_maxunpool_export_without_output_shape_cpu')
-    backend_test.exclude(r'test_mod_mixed_sign_int32_cpu')
-    backend_test.exclude(r'test_mod_mixed_sign_int8_cpu')
     backend_test.exclude(r'test_qlinearmatmul_2D_cpu')
     backend_test.exclude(r'test_qlinearmatmul_3D_cpu')
     backend_test.exclude(r'test_range_float_type_positive_delta_expanded_cpu')