From: "Patrick Lauer" <patrick@gentoo.org>
To: gentoo-commits@lists.gentoo.org
Subject: [gentoo-commits] repo/gentoo:master commit in: dev-libs/rocm-device-libs/files/, dev-libs/rocm-device-libs/
Date: Sun, 28 Jul 2024 08:00:41 +0000 (UTC) [thread overview]
Message-ID: <1722153639.365f43025d18775f1d70b9ef6e30a2403cf80fe6.patrick@gentoo> (raw)
commit: 365f43025d18775f1d70b9ef6e30a2403cf80fe6
Author: Patrick Lauer <patrick <AT> gentoo <DOT> org>
AuthorDate: Sun Jul 28 07:57:38 2024 +0000
Commit: Patrick Lauer <patrick <AT> gentoo <DOT> org>
CommitDate: Sun Jul 28 08:00:39 2024 +0000
URL: https://gitweb.gentoo.org/repo/gentoo.git/commit/?id=365f4302
dev-libs/rocm-device-libs: add 6.1.2
Signed-off-by: Patrick Lauer <patrick <AT> gentoo.org>
dev-libs/rocm-device-libs/Manifest | 1 +
.../files/rocm-device-libs-6.1.2-fix-build.patch | 152 +++++++++++++++++++++
.../rocm-device-libs/rocm-device-libs-6.1.2.ebuild | 84 ++++++++++++
3 files changed, 237 insertions(+)
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}}"
+}
next reply other threads:[~2024-07-28 8:00 UTC|newest]
Thread overview: 3+ messages / expand[flat|nested] mbox.gz Atom feed top
2024-07-28 8:00 Patrick Lauer [this message]
-- strict thread matches above, loose matches on Subject: below --
2023-09-03 6:31 [gentoo-commits] repo/gentoo:master commit in: dev-libs/rocm-device-libs/files/, dev-libs/rocm-device-libs/ 罗百科
2022-08-06 14:23 Benda XU
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=1722153639.365f43025d18775f1d70b9ef6e30a2403cf80fe6.patrick@gentoo \
--to=patrick@gentoo.org \
--cc=gentoo-commits@lists.gentoo.org \
--cc=gentoo-dev@lists.gentoo.org \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox