sci-libs/miopen: add 5.0.2

Changes: patch the sources a bit to support gfx1031

Closes: https://github.com/gentoo/gentoo/pull/25851
Closes: https://bugs.gentoo.org/851102
Signed-off-by: Yiyang Wu <xgreenlandforwyy@gmail.com>
Signed-off-by: Benda Xu <heroxbd@gentoo.org>

5 files changed, 361 insertions, 0 deletions

diff --git a/sci-libs/miopen/Manifest b/sci-libs/miopen/Manifest
index ae3770f95d24..2a8dbb55fead 100644
--- a/sci-libs/miopen/Manifest
+++ b/sci-libs/miopen/Manifest
@@ -1 +1,2 @@
 DIST MIOpen-4.3.0.tar.gz 59405686 BLAKE2B fe91bd91a74023866883d6d0a2a8071a8fe40c4cff2fb4ef58fc6e343a05ac2a731f73e657f4d183ade4e5b7c1fbbe41f3f032918f6e50cb713073aee8d97dc5 SHA512 a8615b7738acfbc7f68d9417b0746c62630d2b48fb8485fafba4db65a4b277a8230f601d249d7e54f89ba25c14176429ca76ad8365a437b09d836b031b0c4fbb
+DIST MIOpen-5.0.2.tar.gz 76294827 BLAKE2B 7b2a1f0e675793aee4a0fa2a270caac8332cda36c8f04cee483cc2882ed987b6e676e9c24a1acf4976a16a10f922b1a6263470b419aa88a29cfcb6d6d4b4cc29 SHA512 a581b45220797904db3e4dd3840f2ef96085f00baf8187c5ab574325a66da4f599dee6496457bb1cc32825b57a13fb0ef35a2ef1bd2a5f449c7e7b9fa64b27d1
diff --git a/sci-libs/miopen/files/miopen-5.0.2-gfx1031.patch b/sci-libs/miopen/files/miopen-5.0.2-gfx1031.patch
new file mode 100644
index 000000000000..15ac67bd3cef
--- /dev/null
+++ b/sci-libs/miopen/files/miopen-5.0.2-gfx1031.patch
@@ -0,0 +1,241 @@
+Index: MIOpen-rocm-5.0.2/src/include/miopen/solver/implicitgemm_util.hpp
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/src/include/miopen/solver/implicitgemm_util.hpp
++++ MIOpen-rocm-5.0.2/src/include/miopen/solver/implicitgemm_util.hpp
+@@ -478,7 +478,7 @@ static inline bool is_use_amd_buffer_loa
+ {
+ #if WORKAROUND_MIOPEN_ISSUE_557
+     const auto device_name = ctx.GetStream().GetDeviceName();
+-    return !StartsWith(device_name, "gfx1030");
++    return !StartsWith(device_name, "gfx1030") && !StartsWith(device_name, "gfx1031");
+ #else
+     return true;
+ #endif
+@@ -487,7 +487,7 @@ static inline bool is_use_amd_buffer_loa
+ static inline bool is_use_v_fmac_f32(const ConvolutionContext& ctx)
+ {
+     const auto device_name = ctx.GetStream().GetDeviceName();
+-    return StartsWith(device_name, "gfx1030");
++    return StartsWith(device_name, "gfx1030") || StartsWith(device_name, "gfx1031");
+ }
+ 
+ static inline bool support_amd_buffer_atomic_fadd(const std::string& device_name)
+@@ -608,7 +608,8 @@ static inline bool IsComposableKernelSup
+            StartsWith(c.GetStream().GetDeviceName(), "gfx906") ||
+            StartsWith(c.GetStream().GetDeviceName(), "gfx908") ||
+            StartsWith(c.GetStream().GetDeviceName(), "gfx90a") ||
+-           StartsWith(c.GetStream().GetDeviceName(), "gfx1030");
++           StartsWith(c.GetStream().GetDeviceName(), "gfx1030")||
++           StartsWith(c.GetStream().GetDeviceName(), "gfx1031");
+ }
+ 
+ // greatest common divisor, aka highest common factor
+Index: MIOpen-rocm-5.0.2/src/kernels/batchnorm_functions.h
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/src/kernels/batchnorm_functions.h
++++ MIOpen-rocm-5.0.2/src/kernels/batchnorm_functions.h
+@@ -159,6 +159,10 @@
+ #define MIO_BN_GFX1030 0
+ #endif
+ 
++#ifndef MIO_BN_GFX1031
++#define MIO_BN_GFX1031 0
++#endif
++
+ #define UNUSED __attribute__((__unused__))
+ 
+ #if(MIO_BN_VARIANT != 4)
+Index: MIOpen-rocm-5.0.2/src/kernels/MIOpenBatchNormActivBwdPerAct.cl
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/src/kernels/MIOpenBatchNormActivBwdPerAct.cl
++++ MIOpen-rocm-5.0.2/src/kernels/MIOpenBatchNormActivBwdPerAct.cl
+@@ -34,7 +34,7 @@
+ #endif
+ 
+ #define MIOPEN_USE_AMDGCN 0
+-#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1
++#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 && MIO_BN_GFX1031 != 1
+ #undef MIOPEN_USE_AMDGCN
+ #define MIOPEN_USE_AMDGCN 1
+ #endif
+Index: MIOpen-rocm-5.0.2/src/kernels/MIOpenBatchNormActivBwdSpatial.cl
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/src/kernels/MIOpenBatchNormActivBwdSpatial.cl
++++ MIOpen-rocm-5.0.2/src/kernels/MIOpenBatchNormActivBwdSpatial.cl
+@@ -32,7 +32,7 @@
+ #endif
+ 
+ #define MIOPEN_USE_AMDGCN 0
+-#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1
++#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 && MIO_BN_GFX1031 != 1
+ #undef MIOPEN_USE_AMDGCN
+ #define MIOPEN_USE_AMDGCN 1
+ #endif
+Index: MIOpen-rocm-5.0.2/src/kernels/MIOpenBatchNormActivFwdTrainSpatial.cl
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/src/kernels/MIOpenBatchNormActivFwdTrainSpatial.cl
++++ MIOpen-rocm-5.0.2/src/kernels/MIOpenBatchNormActivFwdTrainSpatial.cl
+@@ -33,7 +33,7 @@
+ #endif
+ 
+ #define MIOPEN_USE_AMDGCN 0
+-#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1
++#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 && MIO_BN_GFX1031 != 1
+ #undef MIOPEN_USE_AMDGCN
+ #define MIOPEN_USE_AMDGCN 1
+ #endif
+Index: MIOpen-rocm-5.0.2/src/kernels/MIOpenBatchNormBwdSpatial.cl
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/src/kernels/MIOpenBatchNormBwdSpatial.cl
++++ MIOpen-rocm-5.0.2/src/kernels/MIOpenBatchNormBwdSpatial.cl
+@@ -33,7 +33,7 @@
+ #endif
+ 
+ #define MIOPEN_USE_AMDGCN 0
+-#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1
++#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 && MIO_BN_GFX1031 != 1
+ #undef MIOPEN_USE_AMDGCN
+ #define MIOPEN_USE_AMDGCN 1
+ #endif
+Index: MIOpen-rocm-5.0.2/src/kernels/MIOpenBatchNormFwdTrainSpatial.cl
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/src/kernels/MIOpenBatchNormFwdTrainSpatial.cl
++++ MIOpen-rocm-5.0.2/src/kernels/MIOpenBatchNormFwdTrainSpatial.cl
+@@ -33,7 +33,7 @@
+ #endif
+ 
+ #define MIOPEN_USE_AMDGCN 0
+-#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1
++#if defined(__AMDGCN__) && MIO_BN_GFX1030 != 1 && MIO_BN_GFX1031 != 1
+ #undef MIOPEN_USE_AMDGCN
+ #define MIOPEN_USE_AMDGCN 1
+ #endif
+Index: MIOpen-rocm-5.0.2/src/md_graph.cpp
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/src/md_graph.cpp
++++ MIOpen-rocm-5.0.2/src/md_graph.cpp
+@@ -738,8 +738,8 @@ void FusionMDGraph::InitConv(FusionMDGra
+ 
+             add_v21_wino("gfx9", {"gfx900", "gfx906", "gfx908", "gfx90a"}, 1);
+             add_v21_wino("gfx9", {"gfx900", "gfx906", "gfx908", "gfx90a"}, 2);
+-            add_v21_wino("gfx10", {"gfx1011", "gfx1012", "gfx1030"}, 1);
+-            add_v21_wino("gfx10", {"gfx1011", "gfx1012", "gfx1030"}, 2);
++            add_v21_wino("gfx10", {"gfx1011", "gfx1012", "gfx1030", "gfx1031"}, 1);
++            add_v21_wino("gfx10", {"gfx1011", "gfx1012", "gfx1030", "gfx1031"}, 2);
+         }
+     }
+ 
+Index: MIOpen-rocm-5.0.2/src/ocl/fusionopbiasbnactivocl.cpp
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/src/ocl/fusionopbiasbnactivocl.cpp
++++ MIOpen-rocm-5.0.2/src/ocl/fusionopbiasbnactivocl.cpp
+@@ -392,7 +392,8 @@ miopenStatus_t BatchNormBwdTrainFusionOp
+            " -DMIO_BN_USESAVED=" + std::to_string(static_cast<int>(true)) +
+            " -DMIO_BN_VARIANT=" + std::to_string(variant) +
+            " -DMIO_BN_CBA_WRITE_INTERMEDIATE=" + std::to_string(0) +
+-           " -DMIO_BN_GFX1030=" + ((handle.GetDeviceName() == "gfx1030") ? "1" : "0");
++           " -DMIO_BN_GFX1030=" + ((handle.GetDeviceName() == "gfx1030") ? "1" : "0") +
++           " -DMIO_BN_GFX1031=" + ((handle.GetDeviceName() == "gfx1031") ? "1" : "0");
+ 
+     compile_config += add;
+     MIOPEN_LOG_I2(add);
+@@ -607,7 +608,8 @@ miopenStatus_t BatchNormFwdTrainFusionOp
+            " -DMIO_SAVE_MEAN_VARIANCE=" + (saveBatchStats ? "1" : "0") +
+            " -DMIO_RUNNING_RESULT=" + ((savePopStats) ? "1" : "0") +
+            " -DMIO_BN_VARIANT=" + std::to_string(variant) +
+-           " -DMIO_BN_GFX1030=" + ((handle.GetDeviceName() == "gfx1030") ? "1" : "0");
++           " -DMIO_BN_GFX1030=" + ((handle.GetDeviceName() == "gfx1030") ? "1" : "0") +
++           " -DMIO_BN_GFX1031=" + ((handle.GetDeviceName() == "gfx1031") ? "1" : "0");
+ 
+     compile_config += add;
+     MIOPEN_LOG_I2(add);
+Index: MIOpen-rocm-5.0.2/src/target_properties.cpp
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/src/target_properties.cpp
++++ MIOpen-rocm-5.0.2/src/target_properties.cpp
+@@ -54,6 +54,7 @@ static std::string GetDeviceNameFromMap(
+         {"Vega10", "gfx900"},
+         {"gfx901", "gfx900"},
+         {"10.3.0 Sienna_Cichlid 18", "gfx1030"},
++        {"10.3.1 Navi_flounder 18", "gfx1031"},
+     };
+ 
+     const char* const p_asciz = miopen::GetStringEnv(MIOPEN_DEBUG_ENFORCE_DEVICE{});
+Index: MIOpen-rocm-5.0.2/test/CMakeLists.txt
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/test/CMakeLists.txt
++++ MIOpen-rocm-5.0.2/test/CMakeLists.txt
+@@ -37,6 +37,7 @@ option( MIOPEN_TEST_GFX908 "Test on MI10
+ option( MIOPEN_TEST_GFX90A "Test on gfx90a" OFF )
+ option( MIOPEN_TEST_VEGA "Test on Vega10/20 (gfx900, gfx906)" OFF )
+ option( MIOPEN_TEST_GFX1030 "Test on Navi21 (gfx1030)" OFF )
++option( MIOPEN_TEST_GFX1031 "Test on Navi21 (gfx1031)" OFF )
+ option( MIOPEN_TEST_GPU_XNACK_ENABLED "Test as if XNACK mode is enabled" OFF )
+ option( MIOPEN_TEST_CONV Off)
+ option( MIOPEN_TEST_DEEPBENCH Off)
+@@ -74,7 +75,7 @@ endif()
+ # Also we do not detect GPU when target GPU for testing is specified explicitly.
+ set(MIOPEN_TEST_GPU_DETECTION_FAILED FALSE)
+ set(MIOPEN_NO_GPU FALSE)
+-if(NOT (MIOPEN_TEST_VEGA OR MIOPEN_TEST_GFX908 OR MIOPEN_TEST_GFX90A OR MIOPEN_TEST_GFX1030 OR MIOPEN_TEST_HIP_NOGPU))
++if(NOT (MIOPEN_TEST_VEGA OR MIOPEN_TEST_GFX908 OR MIOPEN_TEST_GFX90A OR MIOPEN_TEST_GFX1030 OR MIOPEN_TEST_GFX1031 OR MIOPEN_TEST_HIP_NOGPU))
+     find_program(ROCMINFO
+         NAMES rocminfo
+         PATHS
+@@ -96,6 +97,8 @@ if(NOT (MIOPEN_TEST_VEGA OR MIOPEN_TEST_
+         elseif (NOT ROCMINFO_EXIT_STATUS EQUAL 0)
+             message(WARNING "ROCMINFO FAILED, GPU TYPE UNKNOWN. Manually set respective MIOPEN_TEST_GFX* CMake variable to specify target GPU for testing.")
+             set(MIOPEN_TEST_GPU_DETECTION_FAILED TRUE)
++        elseif(ROCMINFO_OUTPUT MATCHES "gfx1031")
++            set(MIOPEN_TEST_GFX1031 ON)
+         elseif(ROCMINFO_OUTPUT MATCHES "gfx1030")
+             set(MIOPEN_TEST_GFX1030 ON)
+         elseif(ROCMINFO_OUTPUT MATCHES "gfx900|gfx906")
+@@ -122,6 +125,7 @@ message(STATUS "MIOPEN_TEST_VEGA ${MIOPE
+ message(STATUS "MIOPEN_TEST_GFX908 ${MIOPEN_TEST_GFX908}")
+ message(STATUS "MIOPEN_TEST_GFX90A ${MIOPEN_TEST_GFX90A}")
+ message(STATUS "MIOPEN_TEST_GFX1030 ${MIOPEN_TEST_GFX1030}")
++message(STATUS "MIOPEN_TEST_GFX1031 ${MIOPEN_TEST_GFX1031}")
+ message(STATUS "MIOPEN_TEST_GPU_XNACK_ENABLED ${MIOPEN_TEST_GPU_XNACK_ENABLED}")
+ message(STATUS "MIOPEN_TEST_GPU_DETECTION_FAILED ${MIOPEN_TEST_GPU_DETECTION_FAILED}")
+ 
+@@ -164,10 +168,10 @@ endmacro()
+ set_var_to_condition(WORKAROUND_ISSUE_1187_DEFAULT MIOPEN_TEST_GFX90A AND MIOPEN_TEST_FLOAT)
+ option( WORKAROUND_ISSUE_1187 "" ${WORKAROUND_ISSUE_1187_DEFAULT})
+ 
+-set_var_to_condition(WORKAROUND_ISSUE_1148_DEFAULT MIOPEN_TEST_GFX1030 AND MIOPEN_TEST_FLOAT)
++set_var_to_condition(WORKAROUND_ISSUE_1148_DEFAULT MIOPEN_TEST_GFX1030 OR MIOPEN_TEST_GFX1031 AND MIOPEN_TEST_FLOAT)
+ option( WORKAROUND_ISSUE_1148 "" ${WORKAROUND_ISSUE_1148_DEFAULT})
+ 
+-set_var_to_condition(WORKAROUND_ISSUE_1334_DEFAULT MIOPEN_TEST_GFX1030 AND MIOPEN_TEST_FLOAT)
++set_var_to_condition(WORKAROUND_ISSUE_1334_DEFAULT MIOPEN_TEST_GFX1030 OR MIOPEN_TEST_GFX1031 AND MIOPEN_TEST_FLOAT)
+ option( WORKAROUND_ISSUE_1334 "" ${WORKAROUND_ISSUE_1334_DEFAULT})
+ 
+ set_var_to_condition(WORKAROUND_ISSUE_1317_DEFAULT MIOPEN_TEST_OPENCL)
+@@ -216,7 +220,7 @@ if (MIOPEN_NO_GPU)
+             test_pooling3d test_perfdb)
+ endif()
+ 
+-if(MIOPEN_TEST_GFX1030)
++if(MIOPEN_TEST_GFX1030 OR MIOPEN_TEST_GFX1031)
+     if(WORKAROUND_ISSUE_1053 AND MIOPEN_TEST_ALL)
+         list(APPEND SKIP_TESTS test_lrn_test)
+     endif()
+@@ -443,7 +447,7 @@ endfunction()
+ #   If nothing is specified, the default value is taken.
+ #   Default: FLOAT_ENABLED HALF_DISABLED BF16_DISABLED INT8_DISABLED
+ #
+-# GPU types: VEGA, GFX908, GFX90A, GFX1030
++# GPU types: VEGA, GFX908, GFX90A, GFX1030, GFX1031
+ #   VEGA tests are intended to be run on gfx900 or gfx906.
+ #   The option can be enabled or disabled by using '_ENABLED' and '_DISABLED' suffix.
+ #   If nothing is specified, the default value is taken.
+@@ -574,7 +578,7 @@ function(add_custom_test NAME)
+         set_tests_properties(${NAME} PROPERTIES RUN_SERIAL On)
+     endif()
+ 
+-    if(  (is_vega_check OR is_gfx908_check OR is_gfx1030_check OR is_gfx90a_check)
++	if(  (is_vega_check OR is_gfx908_check OR is_gfx1030_check OR is_gfx1031_check OR is_gfx90a_check)
+      AND is_full_check
+      AND is_xnack_on_check
+      AND (is_miotensile_check AND is_mlir_check)
diff --git a/sci-libs/miopen/files/miopen-5.0.2-no-strip.patch b/sci-libs/miopen/files/miopen-5.0.2-no-strip.patch
new file mode 100644
index 000000000000..a7be67e99578
--- /dev/null
+++ b/sci-libs/miopen/files/miopen-5.0.2-no-strip.patch
@@ -0,0 +1,18 @@
+Don't strip for release. Let portage handle stripping.
+Index: MIOpen-rocm-5.0.2/CMakeLists.txt
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/CMakeLists.txt
++++ MIOpen-rocm-5.0.2/CMakeLists.txt
+@@ -78,12 +78,6 @@ option( BUILD_DEV "Build for development
+ option(MIOPEN_ENABLE_FIN "Enable the fin driver for MIOpen"  OFF)
+ 
+ 
+-# Strip symbols for release
+-if(NOT WIN32 AND NOT APPLE)
+-    set(CMAKE_C_FLAGS_RELEASE "${CMAKE_C_FLAGS_RELEASE} -s")
+-    set(CMAKE_CXX_FLAGS_RELEASE "${CMAKE_CXX_FLAGS_RELEASE} -s")
+-endif()
+-
+ rocm_setup_version(VERSION 2.15.0)
+ 
+ list( APPEND CMAKE_MODULE_PATH ${PROJECT_SOURCE_DIR}/cmake )
diff --git a/sci-libs/miopen/files/miopen-5.0.2-strip-xnack-in-flags.patch b/sci-libs/miopen/files/miopen-5.0.2-strip-xnack-in-flags.patch
new file mode 100644
index 000000000000..14d2c5b9daa5
--- /dev/null
+++ b/sci-libs/miopen/files/miopen-5.0.2-strip-xnack-in-flags.patch
@@ -0,0 +1,20 @@
+if options like :xnack- exists in ${AMDGPU_TARGETS}, CMakeLists cannot handle HIP_COMPILER_FLAGS well
+
+the original regex replace should include :+- so xnack- is stripped as well. Otherwise clang complation at MIOpen runtime will fail.
+
+Signed-off-by: Yiyang Wu <xgreenlandforwyy@gmail.com>
+Index: MIOpen-rocm-5.0.2/CMakeLists.txt
+===================================================================
+--- MIOpen-rocm-5.0.2.orig/CMakeLists.txt
++++ MIOpen-rocm-5.0.2/CMakeLists.txt
+@@ -198,8 +198,8 @@ find_package(hip REQUIRED PATHS /opt/roc
+ message(STATUS "Build with HIP ${hip_VERSION}")
+ target_flags(HIP_COMPILER_FLAGS hip::device)
+ # Remove cuda arch flags
+-string(REGEX REPLACE --cuda-gpu-arch=[a-z0-9]+ "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
+-string(REGEX REPLACE --offload-arch=[a-z0-9]+ "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
++string(REGEX REPLACE --cuda-gpu-arch=[a-z0-9:+-]+ "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
++string(REGEX REPLACE --offload-arch=[a-z0-9:+-]+ "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
+ string(REPLACE "$<LINK_LANGUAGE:CXX>" "1" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
+ string(REPLACE "SHELL:" "" HIP_COMPILER_FLAGS "${HIP_COMPILER_FLAGS}")
+ 
diff --git a/sci-libs/miopen/miopen-5.0.2.ebuild b/sci-libs/miopen/miopen-5.0.2.ebuild
new file mode 100644
index 000000000000..1a75457baec2
--- /dev/null
+++ b/sci-libs/miopen/miopen-5.0.2.ebuild
@@ -0,0 +1,81 @@
+# Copyright 1999-2022 Gentoo Authors
+# Distributed under the terms of the GNU General Public License v2
+
+EAPI=8
+
+inherit cmake flag-o-matic
+
+DESCRIPTION="AMD's Machine Intelligence Library"
+HOMEPAGE="https://github.com/ROCmSoftwarePlatform/MIOpen"
+SRC_URI="https://github.com/ROCmSoftwarePlatform/MIOpen/archive/rocm-${PV}.tar.gz -> MIOpen-${PV}.tar.gz"
+
+LICENSE="MIT"
+KEYWORDS="~amd64"
+SLOT="0/$(ver_cut 1-2)"
+
+IUSE="debug test"
+RESTRICT="!test? ( test )"
+
+RDEPEND="
+	dev-util/hip
+	>=dev-db/sqlite-3.17
+	dev-libs/ocl-icd
+	dev-util/rocm-clang-ocl:${SLOT}
+	sci-libs/rocBLAS:${SLOT}
+	>=dev-libs/boost-1.72
+"
+
+DEPEND="${RDEPEND}"
+
+BDEPEND="dev-libs/half:0/1"
+
+S="${WORKDIR}/MIOpen-rocm-${PV}"
+
+PATCHES=(
+	"${FILESDIR}/${PN}-4.2.0-disable-no-inline-boost.patch"
+	"${FILESDIR}/${PN}-4.2.0-gcc11-numeric_limits.patch"
+	"${FILESDIR}/${PN}-5.0.2-strip-xnack-in-flags.patch"
+	"${FILESDIR}/${PN}-4.3.0-fix-interface-include-in-HIP_COMPILER_FLAGS.patch"
+	"${FILESDIR}/${PN}-4.3.0-enable-test.patch"
+	"${FILESDIR}/${PN}-5.0.2-no-strip.patch"
+	"${FILESDIR}/${PN}-5.0.2-gfx1031.patch"
+)
+
+src_prepare() {
+	sed -e "s:/opt/rocm/llvm:""${EPREFIX}""/usr/lib/llvm/roc/ NO_DEFAULT_PATH:" \
+		-e "s:/opt/rocm/hip:""${EPREFIX}""/usr/lib/hip/ NO_DEFAULT_PATH:" \
+		-e '/set( MIOPEN_INSTALL_DIR/s:miopen:${CMAKE_INSTALL_PREFIX}:' \
+		-e '/MIOPEN_TIDY_ERRORS ALL/d' \
+		-i CMakeLists.txt || die
+
+	sed -e "/rocm_install_symlink_subdir(\${MIOPEN_INSTALL_DIR})/d" -i src/CMakeLists.txt || die
+	sed -e "/add_test/s:--build \${CMAKE_CURRENT_BINARY_DIR}:--build ${BUILD_DIR}:" -i test/CMakeLists.txt || die
+
+	sed -e "s:\${AMD_DEVICE_LIBS_PREFIX}/lib:${EPREFIX}/usr/lib/amdgcn/bitcode:" -i cmake/hip-config.cmake || die
+
+	cmake_src_prepare
+}
+
+src_configure() {
+	if ! use debug; then
+		append-cflags "-DNDEBUG"
+		append-cxxflags "-DNDEBUG"
+		CMAKE_BUILD_TYPE="Release"
+	else
+		CMAKE_BUILD_TYPE="Debug"
+	fi
+
+	export CXX="${EPREFIX}/usr/lib/llvm/roc/bin/clang++"
+
+	local mycmakeargs=(
+		-DCMAKE_SKIP_RPATH=ON
+		-DCMAKE_INSTALL_PREFIX="${EPREFIX}/usr"
+		-DMIOPEN_BACKEND=HIP
+		-DBoost_USE_STATIC_LIBS=OFF
+		-DBUILD_TESTS=$(usex test ON OFF)
+		-DMIOPEN_TEST_ALL=$(usex test ON OFF)
+		${AMDGPU_TARGETS+-DAMDGPU_TARGETS="${AMDGPU_TARGETS}"}
+	)
+
+	cmake_src_configure
+}