dev-libs/rocm-device-libs: add 6.1.1

Update remote-id in metadata.

Upstream SRC_URI changed, from separate repo to subproject of ROCm
forked llvm. Borrow unpack function of llvm.org.eclass to unpack the
amd/device-libs component.

Remove shadowed RESTRICT="strip", because .bc files won't be stripped
so there's no need for RESTRICT=strip.

Some tests failed for older arch, probably due to inconsistency between
AMD forked llvm and upstream llvm:

  6 - compile_frexp__gfx600 (Failed)
  7 - compile_fract__gfx600 (Failed)
 12 - compile_fract__gfx700 (Failed)
 17 - compile_fract__gfx803 (Failed)

Signed-off-by: Yiyang Wu <xgreenlandforwyy@gmail.com>
Signed-off-by: Sam James <sam@gentoo.org>

5 files changed, 209 insertions, 0 deletions

diff --git a/dev-libs/rocm-device-libs/Manifest b/dev-libs/rocm-device-libs/Manifest
index 8002ba47f4ad..3e1762b09f72 100644
--- a/dev-libs/rocm-device-libs/Manifest
+++ b/dev-libs/rocm-device-libs/Manifest
@@ -1,3 +1,4 @@
+DIST llvm-project-rocm-6.1.1.tar.gz 196027084 BLAKE2B 3ef0b6e3c47c66fd80289373e6ff8aaff44751f9b380addfae73a18dc388093c0535f230b0cc7528724bc43f6992e2ae6decd3d0d3c700893ca95a6166b7b8dc SHA512 e320d4eeaa6f61ed1cdbf653d67fe887d3ce9dc0d6743b4713502e1cb5318ab8afbe1ee71f8cba07635c54ce532df6683de40ade0e5be4a52e50ce25a9b70818
 DIST rocm-device-libs-5.1.3.tar.gz 242862 BLAKE2B 68d66de897f461e9f876de5fe2214803d4c00665651dea6af0952f0ce579c6704a5ec41b08971fa613ade309a0a85cb611b56b592dc2a25e247183e634ea3378 SHA512 cc3dfb8d4b4841ba777355c537175259d0019159ff462358320674b85082cccd99f6462f60fee66228ddfb88fade043445c1bac62504aa1462ba61b7e2751de7
 DIST rocm-device-libs-5.3.3.tar.gz 245690 BLAKE2B 475c0d818b8b0f090a8daeca2910cd4002e4cdf505d020327f46eb5f864a26937a6a3dfe4ff7b188ebda0f936b1c396f2163bb27b9e2a62c5976e60fa60856ac SHA512 8f6f2fc1534e348e02ba30a25cfc6017f8eab768968b5d0344a5ea7d65c4f0a874072f9e53919c74545814330602ef7c190753c7ff019137230e02f58a5d3a5d
 DIST rocm-device-libs-5.4.3.tar.gz 246095 BLAKE2B eb749346c96d465a5f22831968ccbd71f02749e6aa0d9c2becc0f378641ca0f65c1a131bfd3ed226f838b4208091fcc920b1e31b427adbd69a42881898668e6a SHA512 67b904363a3cff6c15bbd032cbc72cb5cd5f82acaa68c74391dbcf415266e8f35486a496b69b69e1fc0721e0e4e21fb6a6b9c180a46cb59cdcf53916be846ca4
diff --git a/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-llvm-link.patch b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-llvm-link.patch
new file mode 100644
index 000000000000..7a08dc4a1d0d
--- /dev/null
+++ b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-llvm-link.patch
@@ -0,0 +1,28 @@
+https://github.com/ROCm/llvm-project/pull/68
+From 1c7e7f872980a5b15fb3d85f8780e78ce3b715b1 Mon Sep 17 00:00:00 2001
+From: Selene <lixueying@mail.bnu.edu.cn>
+Date: Wed, 1 May 2024 17:38:33 +0800
+Subject: [PATCH] Allow link to llvm shared library for current distros
+
+---
+ amd/device-libs/utils/prepare-builtins/CMakeLists.txt | 6 +++++-
+ 1 file changed, 5 insertions(+), 1 deletion(-)
+
+diff --git a/amd/device-libs/utils/prepare-builtins/CMakeLists.txt b/amd/device-libs/utils/prepare-builtins/CMakeLists.txt
+index 63661962860a..079dc08e3419 100644
+--- a/utils/prepare-builtins/CMakeLists.txt
++++ b/utils/prepare-builtins/CMakeLists.txt
+@@ -26,5 +26,9 @@ set_target_properties(prepare-builtins PROPERTIES
+   CXX_STANDARD_REQUIRED Yes
+   CXX_EXTENSIONS No)
+ llvm_update_compile_flags(prepare-builtins)
+-llvm_map_components_to_libnames(llvm_libs support core bitreader bitwriter)
++if (LLVM_LINK_LLVM_DYLIB)
++  set(llvm_libs LLVM)
++else()
++  llvm_map_components_to_libnames(llvm_libs support core bitreader bitwriter)
++endif()
+ target_link_libraries(prepare-builtins ${llvm_libs})
+-- 
+2.44.0
+
diff --git a/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures.patch b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures.patch
new file mode 100644
index 000000000000..856b5cf7e392
--- /dev/null
+++ b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures.patch
@@ -0,0 +1,64 @@
+Modified from https://github.com/ROCm/llvm-project/commit/7c2188cbc193f2b4dd5394f17404b44340001f30.patch
+From 7c2188cbc193f2b4dd5394f17404b44340001f30 Mon Sep 17 00:00:00 2001
+From: Matt Arsenault <Matthew.Arsenault@amd.com>
+Date: Thu, 4 Jan 2024 19:34:47 +0700
+Subject: [PATCH] device-libs: Fix input file path test failures
+
+The test file input paths broke with the move to the mono-repo. Some of
+the constant folding values are host dependent, so update the values to
+what works for me now. Not really sure what else I can do about these.
+
+Change-Id: Ic764f637bb9532fcede9bfb9ce3886a2b7d467d8
+---
+ amd/device-libs/test/compile/frexp.cl       | 10 ++++++----
+ amd/device-libs/test/compile/lgamma_r.cl    |  2 +-
+ 3 files changed, 8 insertions(+), 9 deletions(-)
+
+diff --git a/test/compile/frexp.cl b/test/compile/frexp.cl
+index 780c541500021b..b3181ce87db711 100644
+--- a/test/compile/frexp.cl
++++ b/test/compile/frexp.cl
+@@ -5,11 +5,11 @@
+ // later.
+ 
+ // GCN-LABEL: {{^}}test_frexp_f32:
+-// GFX600-DAG: s_movk_i32 [[CLASS_MASK:s[0-9]+]], 0x1f8
++// GFX600-DAG: s_mov_b32 [[INF:s[0-9]+]], 0x7f80000
+ // GFX600-DAG: v_frexp_mant_f32{{(_e32)?}} [[MANT:v[0-9]+]], [[SRC:v[0-9]+]]
+ // GFX600-DAG: v_frexp_exp_i32_f32{{(_e32)?}} [[EXP:v[0-9]+]], [[SRC:v[0-9]+]]
+ 
+-// GFX600-DAG: v_cmp_class_f32{{(_e64)?}} [[CMP:(vcc|s{{\[[0-9]+:[0-9]+\]}})]], [[SRC]], [[CLASS_MASK]]
++// GFX600-DAG: v_cmp_lt_f32{{(_e64)?}} [[CMP:(vcc|s{{\[[0-9]+:[0-9]+\]}})]], |[[SRC]]|, [[INF]]
+ 
+ // GFX600-DAG: v_cndmask_b32{{(_e32)?|(e64)?}} v{{[0-9]+}}, [[SRC]], [[MANT]], [[CMP]]
+ // GFX600-DAG: v_cndmask_b32{{(_e32)?|(e64)?}} v{{[0-9]+}}, 0, [[EXP]], [[CMP]]
+@@ -30,12 +30,14 @@ kernel void test_frexp_f32(global float* restrict out0,
+ }
+ 
+ // GCN-LABEL: {{^}}test_frexp_f64:
++// GFX600: s_mov_b32 s{{[0-9]+}}, 0{{$}}
+ 
+-// GFX600-DAG: s_movk_i32 [[CLASS_MASK:s[0-9]+]], 0x1f8
++// GFX600-DAG: s_mov_b32 s[[INF_LO:[0-9]+]], 0{{$}}
++// GFX600-DAG: s_mov_b32 s[[INF_HI:[0-9]+]], 0x7ff00000{{$}}
+ // GFX600-DAG: v_frexp_mant_f64{{(_e32)?}} v{{\[}}[[MANT_LO:[0-9]+]]:[[MANT_HI:[0-9]+]]{{\]}}, [[SRC:v\[[0-9]+:[0-9]+\]]]
+ // GFX600-DAG: v_frexp_exp_i32_f64{{(_e32)?}} [[EXP:v[0-9]+]], [[SRC:v\[[0-9]+:[0-9]+\]]]
+ 
+-// GFX600-DAG: v_cmp_class_f64{{(_e64)?}} [[CMP:(vcc|s{{\[[0-9]+:[0-9]+\]}})]], [[SRC]], [[CLASS_MASK]]
++// GFX600-DAG: v_cmp_lt_f64{{(_e64)?}} [[CMP:(vcc|s{{\[[0-9]+:[0-9]+\]}})]], |[[SRC]]|, s{{\[}}[[INF_LO]]:[[INF_HI]]{{\]}}
+ 
+ // GFX600-DAG: v_cndmask_b32{{(_e32)?|(e64)?}} v{{[0-9]+}}, v{{[0-9]+}}, v[[MANT_LO]], [[CMP]]
+ // GFX600-DAG: v_cndmask_b32{{(_e32)?|(e64)?}} v{{[0-9]+}}, v{{[0-9]+}}, v[[MANT_HI]], [[CMP]]
+diff --git a/test/compile/lgamma_r.cl b/test/compile/lgamma_r.cl
+index 1e1984226cd55c..56d1ba15f761f2 100644
+--- a/test/compile/lgamma_r.cl
++++ b/test/compile/lgamma_r.cl
+@@ -66,7 +66,7 @@ kernel void constant_fold_lgamma_r_f32(volatile global float* out,
+     out[0] = test_lgamma_r(0x1.0p-127f, sign_out);
+ 
+     // CONSTANTFOLD-NEXT: store volatile i32 1,
+-    // CONSTANTFOLD-NEXT: store volatile float 0x419DE28040000000,
++    // CONSTANTFOLD-NEXT: store volatile float 0x419DE28060000000,
+     out[0] = test_lgamma_r(nextafter(0x1.0p+23f, __builtin_inff()), sign_out);
+ 
+     // CONSTANTFOLD-NEXT: store volatile i32 1,
diff --git a/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures2.patch b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures2.patch
new file mode 100644
index 000000000000..c0559e9468f6
--- /dev/null
+++ b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.0-fix-test-failures2.patch
@@ -0,0 +1,44 @@
+Modified from https://github.com/ROCm/llvm-project/commit/794ebeffcafbf6f4d86cb1bfd7a5a0d1d30f1fc7
+From 794ebeffcafbf6f4d86cb1bfd7a5a0d1d30f1fc7 Mon Sep 17 00:00:00 2001
+From: Matt Arsenault <Matthew.Arsenault@amd.com>
+Date: Thu, 4 Jan 2024 20:40:30 +0700
+Subject: [PATCH] device-libs: Fix native_rsqrt test failures for f16-as-f32
+ case
+
+At some point the code quality regressed for these.
+
+Change-Id: Ib5c4c1d093a6056dd6213ed6eaf157ad806101fb
+---
+ amd/device-libs/test/compile/native_rsqrt.cl | 13 +++++++++----
+ 1 file changed, 9 insertions(+), 4 deletions(-)
+
+diff --git a/test/compile/native_rsqrt.cl b/test/compile/native_rsqrt.cl
+index 94ef375e4d522f..291cc6f32c3351 100644
+--- a/test/compile/native_rsqrt.cl
++++ b/test/compile/native_rsqrt.cl
+@@ -6,16 +6,21 @@
+ 
+ half __ocml_native_rsqrt_f16(half);
+ 
+-// GCN: {{^}}test_native_rsqrt_f16:
+-// GFX600: v_rsq_f32
+-// GFX700: v_rsq_f32
++// FIXME: Promoted case using full expansion
++// GCN-LABEL: {{^}}test_native_rsqrt_f16:
++// GFX600: v_sqrt_f32
++// GFX600: v_rcp_f32
++
++// GFX700: v_sqrt_f32
++// GFX700: v_rcp_f32
++
+ // GFX803: v_rsq_f16
+ kernel void test_native_rsqrt_f16(global half* restrict out, global half* restrict in) {
+     int id = get_local_id(0);
+     out[id] = __ocml_native_rsqrt_f16(in[id]);
+ }
+ 
+-// GCN: {{^}}test_native_rsqrt_f32:
++// GCN-LABEL: {{^}}test_native_rsqrt_f32:
+ // GCN: v_rsq_f32
+ kernel void test_native_rsqrt_f32(global float* restrict out, global float* restrict in) {
+     int id = get_local_id(0);
diff --git a/dev-libs/rocm-device-libs/rocm-device-libs-6.1.1.ebuild b/dev-libs/rocm-device-libs/rocm-device-libs-6.1.1.ebuild
new file mode 100644
index 000000000000..a4bd9f6a52e9
--- /dev/null
+++ b/dev-libs/rocm-device-libs/rocm-device-libs-6.1.1.ebuild
@@ -0,0 +1,72 @@
+# Copyright 1999-2024 Gentoo Authors
+# Distributed under the terms of the GNU General Public License v2
+
+EAPI=8
+
+LLVM_COMPAT=( 18 )
+inherit cmake llvm-r1
+
+MY_P=llvm-project-rocm-${PV}
+components=( "amd/device-libs" )
+
+if [[ ${PV} == *9999 ]] ; then
+	EGIT_REPO_URI="https://github.com/ROCm/llvm-project"
+	inherit git-r3
+	S="${WORKDIR}/${P}/${components[0]}"
+else
+	SRC_URI="https://github.com/ROCm/llvm-project/archive/rocm-${PV}.tar.gz -> ${MY_P}.tar.gz"
+	S="${WORKDIR}/${MY_P}/${components[0]}"
+	KEYWORDS="~amd64"
+fi
+
+DESCRIPTION="Radeon Open Compute Device Libraries"
+HOMEPAGE="https://github.com/ROCm/ROCm-Device-Libs"
+
+LICENSE="MIT"
+SLOT="0/$(ver_cut 1-2)"
+IUSE="test"
+RESTRICT="!test? ( test )"
+
+RDEPEND="
+	$(llvm_gen_dep '
+		sys-devel/clang:${LLVM_SLOT}
+	')
+"
+DEPEND="${RDEPEND}"
+
+CMAKE_BUILD_TYPE=Release
+
+PATCHES=(
+	"${FILESDIR}/${PN}-5.5.0-test-bitcode-dir.patch"
+	"${FILESDIR}/${PN}-6.1.0-fix-llvm-link.patch"
+	"${FILESDIR}/${PN}-6.0.0-add-gws-attribute.patch"
+	"${FILESDIR}/${PN}-6.1.0-fix-test-failures.patch"
+	"${FILESDIR}/${PN}-6.1.0-fix-test-failures2.patch"
+	)
+
+src_unpack() {
+	if [[ ${PV} == *9999 ]] ; then
+		git-r3_fetch
+		git-r3_checkout '' . '' "${components[@]}"
+	else
+		archive="${MY_P}.tar.gz"
+		ebegin "Unpacking from ${archive}"
+		tar -x -z -o \
+			-f "${DISTDIR}/${archive}" \
+			"${components[@]/#/${MY_P}/}" || die
+		eend ${?}
+	fi
+}
+
+src_prepare() {
+	sed -e "s:amdgcn/bitcode:lib/amdgcn/bitcode:" -i "${S}/cmake/OCL.cmake" || die
+	sed -e "s:amdgcn/bitcode:lib/amdgcn/bitcode:" -i "${S}/cmake/Packages.cmake" || die
+	cmake_src_prepare
+}
+
+src_configure() {
+	local mycmakeargs=(
+		-DLLVM_DIR="$(get_llvm_prefix)"
+	)
+	cmake_src_configure
+}