From ca6512908d44cd56c622b523dc59cb5e79cd39a1 Mon Sep 17 00:00:00 2001 From: Zakor Gyula <126694206+gyulaz-htec@users.noreply.github.com> Date: Mon, 27 May 2024 21:01:17 +0200 Subject: [PATCH 1/2] Fix mod compile issue on GPU (#2268) (#3086) --- src/include/migraphx/op/mod.hpp | 3 +-- .../kernels/include/migraphx/kernels/math.hpp | 18 +++++++++++++++++- test/py/onnx_backend_test.py | 2 -- 3 files changed, 18 insertions(+), 5 deletions(-) diff --git a/src/include/migraphx/op/mod.hpp b/src/include/migraphx/op/mod.hpp index f1a48e3c58f..38f947a3587 100644 --- 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 { 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 index 5a6cca7bc24..da00ff9c781 100644 --- 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>(a, b); } +template ())> +constexpr T mod(const T& a, const T& b) +{ + if constexpr(is_integral{}) + // onnx mod operator requires numpy style modulus + return ((a % b) + b) % b; + return static_cast(fmod(remainder(a, b) + b, b)); +} + +template {} and not is_any_vec())> +constexpr auto mod(const T& a, const U& b) +{ + return mod>(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 index 353bcea3944..2d847c97300 100644 --- 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') From 2dcca47322b967fc99093ec74d940d261d2713a4 Mon Sep 17 00:00:00 2001 From: Chris Austen Date: Mon, 27 May 2024 15:03:01 -0400 Subject: [PATCH 2/2] Move up pybind to support new builds of ONNX Runtime (#3121) --- requirements.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/requirements.txt b/requirements.txt index 49b42573add..ff97144acf2 100755 --- 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