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

Signed-off-by: Patrick Lauer <patrick@gentoo.org>

3 files changed, 237 insertions, 0 deletions

diff --git a/dev-libs/rocm-device-libs/Manifest b/dev-libs/rocm-device-libs/Manifest
index a3682ad63840..91369c2cd3ee 100644
--- a/dev-libs/rocm-device-libs/Manifest
+++ b/dev-libs/rocm-device-libs/Manifest
@@ -1,4 +1,5 @@
 DIST llvm-project-rocm-6.1.1.tar.gz 196027084 BLAKE2B 3ef0b6e3c47c66fd80289373e6ff8aaff44751f9b380addfae73a18dc388093c0535f230b0cc7528724bc43f6992e2ae6decd3d0d3c700893ca95a6166b7b8dc SHA512 e320d4eeaa6f61ed1cdbf653d67fe887d3ce9dc0d6743b4713502e1cb5318ab8afbe1ee71f8cba07635c54ce532df6683de40ade0e5be4a52e50ce25a9b70818
+DIST llvm-project-rocm-6.1.2.tar.gz 195992927 BLAKE2B d821f29f2f1f7c1ff414c63a710281f16d2a394b21f3365d01b86710cc09ed27e514b49fb744bf6a36b38815afa56cc26d44f0238f38479a0c2db9bf9989f389 SHA512 5f7e5dbe5976141de35e96e603624bd9d5a2c08b0690ba9fcd81d1b32f540f94bb9f4b74539e2838fd60ae1312dbe5e0b429ba80a03871782cdf3bd834940ce0
 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.2-fix-build.patch b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.2-fix-build.patch
new file mode 100644
index 000000000000..e717627893b9
--- /dev/null
+++ b/dev-libs/rocm-device-libs/files/rocm-device-libs-6.1.2-fix-build.patch
@@ -0,0 +1,152 @@
+commit 83ef5b48800a47cc30b3d4bcfdf31de9c3bd0dc8
+Author: Patrick Lauer <patrick@gentoo.org>
+Date:   Sun Jul 28 07:43:54 2024 +0000
+
+    Revert "ockl: Don't use wave32 ballot builtin"
+    
+    This reverts commit 066a0b2716b7ade96a2b3e79e5ddcd0c110e9f98.
+
+diff --git a/ockl/src/dm.cl b/ockl/src/dm.cl
+index 18efc54203b7..a3f06c448aee 100644
+--- a/ockl/src/dm.cl
++++ b/ockl/src/dm.cl
+@@ -287,14 +287,9 @@ first(__global void * v)
+     return __builtin_astype(w2, __global void *);
+ }
+ 
+-// Read val from one active lane whose predicate is one.
+-// If no lanes have the predicate set, return none
+-// This is like first, except that first may not have its predicate set
++REQUIRES_WAVE64
+ static uint
+-elect_uint(int pred, uint val, uint none)
+-{
+-  // Pretend wave32 doesn't exist. The wave64 ballot works, and the high half
+-  // will fold out as 0.
++elect_uint_wave64(int pred, uint val, uint none) {
+     uint ret = none;
+ 
+     ulong mask = __builtin_amdgcn_ballot_w64(pred != 0);
+@@ -306,14 +301,51 @@ elect_uint(int pred, uint val, uint none)
+     return ret;
+ }
+ 
+-// Count the number of nonzero arguments across the wave
++REQUIRES_WAVE32
+ static uint
+-votes(bool b)
++elect_uint_wave32(int pred, uint val, uint none) {
++    uint ret = none;
++    uint mask = __builtin_amdgcn_ballot_w32(pred != 0);
++    if (mask != 0U) {
++        uint l = __ockl_ctz_u32(mask);
++        ret = __builtin_amdgcn_ds_bpermute(l << 2, val);
++    }
++
++    return ret;
++}
++
++// Read val from one active lane whose predicate is one.
++// If no lanes have the predicate set, return none
++// This is like first, except that first may not have its predicate set
++static uint
++elect_uint(int pred, uint val, uint none)
++{
++    return __oclc_wavefrontsize64 ?  elect_uint_wave64(pred, val, none) : elect_uint_wave32(pred, val, none);
++}
++
++REQUIRES_WAVE64
++static uint
++votes_wave64(bool b)
+ {
+     ulong mask = __builtin_amdgcn_ballot_w64(b);
+     return __builtin_popcountl(mask);
+ }
+ 
++REQUIRES_WAVE32
++static uint
++votes_wave32(bool b)
++{
++    uint mask = __builtin_amdgcn_ballot_w32(b);
++    return __builtin_popcount(mask);
++}
++
++// Count the number of nonzero arguments across the wave
++static uint
++votes(bool b)
++{
++    return __oclc_wavefrontsize64 ?  votes_wave64(b) : votes_wave32(b);
++}
++
+ // The kind of the smallest block that can hold sz bytes
+ static uint
+ size_to_kind(uint sz)
+diff --git a/ockl/src/wfaas.cl b/ockl/src/wfaas.cl
+index 3861a5bb3eab..4dab97cea5f7 100644
+--- a/ockl/src/wfaas.cl
++++ b/ockl/src/wfaas.cl
+@@ -21,25 +21,60 @@ static int optimizationBarrierHack(int in_val)
+     return out_val;
+ }
+ 
++REQUIRES_WAVE32
++static bool wfany_impl_w32(int e) {
++    return __builtin_amdgcn_ballot_w32(e) != 0;
++}
++
++REQUIRES_WAVE64
++static bool wfany_impl_w64(int e) {
++    return __builtin_amdgcn_ballot_w64(e) != 0;
++}
++
+ ATTR bool
+ OCKL_MANGLE_I32(wfany)(int e)
+ {
+     e = optimizationBarrierHack(e);
+-    return __builtin_amdgcn_ballot_w64(e) != 0;
++    return __oclc_wavefrontsize64 ?
++        wfany_impl_w64(e) : wfany_impl_w32(e);
++}
++
++REQUIRES_WAVE32
++static bool wfall_impl_w32(int e) {
++    return __builtin_amdgcn_ballot_w32(e) == __builtin_amdgcn_read_exec_lo();
++}
++
++REQUIRES_WAVE64
++static bool wfall_impl_w64(int e) {
++    return __builtin_amdgcn_ballot_w64(e) == __builtin_amdgcn_read_exec();
+ }
+ 
+ ATTR bool
+ OCKL_MANGLE_I32(wfall)(int e)
+ {
+     e = optimizationBarrierHack(e);
+-    return __builtin_amdgcn_ballot_w64(e) == __builtin_amdgcn_read_exec();
++    return __oclc_wavefrontsize64 ?
++        wfall_impl_w64(e) : wfall_impl_w32(e);
++}
++
++
++REQUIRES_WAVE32
++static bool wfsame_impl_w32(int e) {
++    uint u = __builtin_amdgcn_ballot_w32(e);
++    return (u == 0) | (u == __builtin_amdgcn_read_exec_lo());
++}
++
++REQUIRES_WAVE64
++static bool wfsame_impl_w64(int e) {
++    ulong u = __builtin_amdgcn_ballot_w64(e);
++    return (u == 0UL) | (u == __builtin_amdgcn_read_exec());
+ }
+ 
+ ATTR bool
+ OCKL_MANGLE_I32(wfsame)(int e)
+ {
+     e = optimizationBarrierHack(e);
+-    ulong u = __builtin_amdgcn_ballot_w64(e);
+-    return (u == 0UL) | (u == __builtin_amdgcn_read_exec());
++    return __oclc_wavefrontsize64 ?
++        wfsame_impl_w64(e) : wfsame_impl_w32(e);
+ }
+ 
diff --git a/dev-libs/rocm-device-libs/rocm-device-libs-6.1.2.ebuild b/dev-libs/rocm-device-libs/rocm-device-libs-6.1.2.ebuild
new file mode 100644
index 000000000000..a6d5dd67438b
--- /dev/null
+++ b/dev-libs/rocm-device-libs/rocm-device-libs-6.1.2.ebuild
@@ -0,0 +1,84 @@
+# 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 )"
+
+BDEPEND="
+	dev-build/rocm-cmake
+	$(llvm_gen_dep '
+		sys-devel/clang:${LLVM_SLOT}
+		sys-devel/lld:${LLVM_SLOT}
+	')
+"
+
+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"
+	"${FILESDIR}/${PN}-6.1.2-fix-build.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)"
+	)
+	# do not trust CMake with autoselecting Clang, as it autoselects the latest one
+	# producing too modern LLVM bitcode and causing linker errors in other packages
+	CC="$(get_llvm_prefix)/bin/clang" CXX="$(get_llvm_prefix)/bin/clang++" cmake_src_configure
+}
+
+src_install() {
+	cmake_src_install
+	local CLANG_EXE="$(get_llvm_prefix)/bin/clang"
+	# install symlink, so that clang won't ask for "--rocm-device-lib-path" flag anymore
+	local bitcodedir="$("${CLANG_EXE}" -print-resource-dir)/$(get_libdir)/amdgcn/bitcode"
+	dosym -r "/usr/lib/amdgcn/bitcode" "${bitcodedir#${EPREFIX}}"
+}