gentoo auto-resync : 22:02:2023 - 20:29:22

3 files changed, 267 insertions, 0 deletions

diff --git a/dev-libs/rocm-device-libs/Manifest b/dev-libs/rocm-device-libs/Manifest
index 4638233aca30..a150a1063115 100644
--- a/dev-libs/rocm-device-libs/Manifest
+++ b/dev-libs/rocm-device-libs/Manifest
@@ -1,7 +1,10 @@
 AUX rocm-device-libs-5.1.3-llvm-link.patch 694 BLAKE2B a67a1207bb23ceaa304f2fd59000a14a8bd232fd6eaf3510370f148b0bd18630ec4309a3d2ee2c5b8045114841a0548e840601abb153fb4adfc1209dd16b0509 SHA512 75f6fb5f3aab7f7951d110b9097a1521adb261f286b5e04867fd2753bbfcc29987440fd929c9dac0ac90fde98616ce7693c2b5041d4d07c437087ee17303a22e
 AUX rocm-device-libs-5.1.3-test-bitcode-dir.patch 1181 BLAKE2B 0358dfc81ed98acd2a2bb399b7d8aa72c2a0bf20da3c45ad02d5e1c46df7f98bc18efe28867725e0fc1df041c1f1147bc25a2a5aa67b18f53d859dbf66aa71c7 SHA512 a0ffdda4ab42978f3b09bab424f91267bfc267f751d489aaee9c6eb3979ae5efd8130129e4b476fdcfbc73b779d9aad79fda60a8f7a64e8a5b35516d8fe6979f
+AUX rocm-device-libs-5.4.3-Revert-Update-counters-for-gfx11.patch 6801 BLAKE2B ae3be5c1d8a7a84e8ddee2471628804365dd1d9b1ae3aa981b5e47015b82c115d5da5ea6fedcacc3ccf11c768af8130375493fdbcc9a53567e633cc57e7e750e SHA512 c28a08253025b150523282231c64e5d310206e6e406efdc2a0f39ce69b2f1663fed6991e27b2124eb3434c51efb8d172eceb752f366bc2aab7b3eb63ad9dfa35
 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
 EBUILD rocm-device-libs-5.1.3-r1.ebuild 1184 BLAKE2B bef07c38844eb218c333cf2fb0014881aff83837092dc62aa29d4971e9782da6a38a7db2dd599f6c8aff0c2bd60c94550d72300d8ee2a6db4ea2688820035803 SHA512 814009470618c7cea132ba28d3aa27eda6a55d65646efa54fc3ce0aa11e247af147f175956d2f7466474d891dd5aef80fe9adcc6bf871352d6f7c20c7f802832
 EBUILD rocm-device-libs-5.3.3.ebuild 1184 BLAKE2B bef07c38844eb218c333cf2fb0014881aff83837092dc62aa29d4971e9782da6a38a7db2dd599f6c8aff0c2bd60c94550d72300d8ee2a6db4ea2688820035803 SHA512 814009470618c7cea132ba28d3aa27eda6a55d65646efa54fc3ce0aa11e247af147f175956d2f7466474d891dd5aef80fe9adcc6bf871352d6f7c20c7f802832
+EBUILD rocm-device-libs-5.4.3.ebuild 1250 BLAKE2B 6013dd693a8120b5eeac654c9bad195428c34059eaef1bd40eaaf856f2edf331d8a3726fa46c077fe0aa62919f3ba68f518cfd9441d23020a06c293cd45f1042 SHA512 df101563920183c6a255f1461e718e28a01b4db355e3807c5a3e7b948627337950105deb7f77a857d489d9591c959909ad1e1957e4f21e3dffba08ee14514469
 MISC metadata.xml 378 BLAKE2B d4defdfb2ce3bab811d9c2ac75b92b7892d46d87dad9f89c26cde140b3b167b43bf4022128dee45e063e332d6dfafc3ee8d78b0112909a7ce2d4a4a44c6cfa14 SHA512 0732d51019f0d27ffb1ffea32d76cad8a6d82927cb8788ff983cdb8dcc472015dbe36ba568bda7459d0b32837e89e1a32064ab7dfcefd2b88458cbb4438f452a
diff --git a/dev-libs/rocm-device-libs/files/rocm-device-libs-5.4.3-Revert-Update-counters-for-gfx11.patch b/dev-libs/rocm-device-libs/files/rocm-device-libs-5.4.3-Revert-Update-counters-for-gfx11.patch
new file mode 100644
index 000000000000..bf9b2c372600
--- /dev/null
+++ b/dev-libs/rocm-device-libs/files/rocm-device-libs-5.4.3-Revert-Update-counters-for-gfx11.patch
@@ -0,0 +1,216 @@
+From 8ce920dddac9846254aaf6261bafd8b22976b04e Mon Sep 17 00:00:00 2001
+From: Jeremy Newton <alexjnewt@hotmail.com>
+Date: Sun, 18 Dec 2022 20:48:21 -0500
+Subject: [PATCH] Revert "Update counters for gfx11"
+
+This reverts commit 85f95b94960c6f7ff4ff0242a399deb4a204fb6a.
+---
+ doc/OCKL.md       |  4 ++--
+ ockl/inc/ockl.h   |  3 ---
+ ockl/src/dm.cl    | 15 +++++++++++----
+ ockl/src/mtime.cl | 35 ++---------------------------------
+ ockl/src/wait.cl  | 18 +++++++++---------
+ 5 files changed, 24 insertions(+), 51 deletions(-)
+
+diff --git a/doc/OCKL.md b/doc/OCKL.md
+index 07574f6..05c5c49 100644
+--- a/doc/OCKL.md
++++ b/doc/OCKL.md
+@@ -99,8 +99,8 @@ The following table lists the available functions along with a brief description
+ | `int __ockl_mul24_i32(int,int);` | Multiply assuming operands fit in 24 bits |
+ | `uint __ockl_mul24_u32(uint,uint);` | |
+ | - | |
+-| `ulong __ockl_cyclectr_u64(void);` | Current value of free running 64-bit clock counter |
+-| `ulong __ockl_steadyctr_u64(void);` | Current value of constant speed 64-bit clock counter |
++| `ulong __ockl_memtime_u64(void);` | Current value of free running 64-bit clock counter |
++| `ulong __ockl_memrealtime_u64(void);` | Current value of constant speed 64-bit clock counter |
+ | - | |
+ | `uint __ockl_activelane_u32(void);` | Index of currently lane counting only active lanes in wavefront |
+ | - | |
+diff --git a/ockl/inc/ockl.h b/ockl/inc/ockl.h
+index d0b98d4..6300279 100644
+--- a/ockl/inc/ockl.h
++++ b/ockl/inc/ockl.h
+@@ -143,9 +143,6 @@ DECL_OCKL_NULLARY_U32(activelane)
+ 
+ DECL_OCKL_NULLARY_U64(memtime)
+ DECL_OCKL_NULLARY_U64(memrealtime)
+-DECL_OCKL_NULLARY_U64(cyclectr)
+-DECL_OCKL_NULLARY_U64(steadyctr)
+-
+ 
+ extern half OCKL_MANGLE_T(wfred_add,f16)(half x);
+ extern float OCKL_MANGLE_T(wfred_add,f32)(float x);
+diff --git a/ockl/src/dm.cl b/ockl/src/dm.cl
+index 245b4a1..26373dd 100644
+--- a/ockl/src/dm.cl
++++ b/ockl/src/dm.cl
+@@ -201,6 +201,13 @@ get_heap_ptr(void) {
+     }
+ }
+ 
++// realtime
++__attribute__((target("s-memrealtime"))) static ulong
++realtime(void)
++{
++    return __builtin_amdgcn_s_memrealtime();
++}
++
+ // The actual number of blocks in a slab with blocks of kind k
+ static uint
+ num_blocks(kind_t k)
+@@ -466,7 +473,7 @@ new_slab_wait(__global heap_t *hp, kind_t k)
+     uint aid = __ockl_activelane_u32();
+     if (aid == 0) {
+         ulong expected = AL(&hp->salloc_time[k].value, memory_order_relaxed);
+-        ulong now = __ockl_steadyctr_u64();
++        ulong now = realtime();
+         ulong dt = now - expected;
+         if  (dt < SLAB_TICKS)
+             __ockl_rtcwait_u32(SLAB_TICKS - (uint)dt);
+@@ -480,7 +487,7 @@ grow_recordable_wait(__global heap_t *hp, kind_t k)
+     uint aid = __ockl_activelane_u32();
+     if (aid == 0) {
+         ulong expected = AL(&hp->grow_time[k].value, memory_order_relaxed);
+-        ulong now = __ockl_steadyctr_u64();
++        ulong now = realtime();
+         ulong dt = now - expected;
+         if  (dt < GROW_TICKS)
+             __ockl_rtcwait_u32(GROW_TICKS - (uint)dt);
+@@ -540,7 +547,7 @@ try_grow_num_recordable_slabs(__global heap_t *hp, kind_t k)
+     uint ret = GROW_BUSY;
+     if (aid == 0) {
+         ulong expected = AL(&hp->grow_time[k].value, memory_order_relaxed);
+-        ulong now = __ockl_steadyctr_u64();
++        ulong now = realtime();
+         if (now - expected >= GROW_TICKS &&
+             ACE(&hp->grow_time[k].value, &expected, now, memory_order_relaxed))
+                 ret = GROW_FAILURE;
+@@ -687,7 +694,7 @@ try_allocate_new_slab(__global heap_t *hp, kind_t k)
+ 
+         if (aid == 0) {
+             ulong expected = AL(&hp->salloc_time[k].value, memory_order_relaxed);
+-            ulong now = __ockl_steadyctr_u64();
++            ulong now = realtime();
+             if (now - expected >= SLAB_TICKS &&
+                 ACE(&hp->salloc_time[k].value, &expected, now, memory_order_relaxed))
+                 ret = (__global sdata_t *)0;
+diff --git a/ockl/src/mtime.cl b/ockl/src/mtime.cl
+index 43f4161..543aaa3 100644
+--- a/ockl/src/mtime.cl
++++ b/ockl/src/mtime.cl
+@@ -5,48 +5,17 @@
+  * License. See LICENSE.TXT for details.
+  *===------------------------------------------------------------------------*/
+ 
+-#include "oclc.h"
+ #include "ockl.h"
+ 
+-__attribute__((target("s-memrealtime"))) static ulong
+-mem_realtime(void)
+-{
+-    return __builtin_amdgcn_s_memrealtime();
+-}
+-
+-__attribute__((target("gfx11-insts"))) static ulong
+-msg_realtime(void)
+-{
+-    return __builtin_amdgcn_s_sendmsg_rtnl(0x83);
+-}
+-
+-// Deprecated
+ __attribute__((target("s-memtime-inst"))) ulong
+ OCKL_MANGLE_U64(memtime)(void)
+ {
+     return __builtin_amdgcn_s_memtime();
+ }
+ 
+-// Deprecated
+-ulong
++__attribute__((target("s-memrealtime"))) ulong
+ OCKL_MANGLE_U64(memrealtime)(void)
+ {
+-    return mem_realtime();
+-}
+-
+-ulong
+-OCKL_MANGLE_U64(cyclectr)(void)
+-{
+-    return __builtin_readcyclecounter();
+-}
+-
+-ulong
+-OCKL_MANGLE_U64(steadyctr)(void)
+-{
+-    if (__oclc_ISA_version >= 11000) {
+-        return msg_realtime();
+-    } else {
+-        return mem_realtime();
+-    }
++    return __builtin_amdgcn_s_memrealtime();
+ }
+ 
+diff --git a/ockl/src/wait.cl b/ockl/src/wait.cl
+index 49b038e..b249599 100644
+--- a/ockl/src/wait.cl
++++ b/ockl/src/wait.cl
+@@ -10,47 +10,47 @@
+ #include "ockl.h"
+ #include "oclc.h"
+ 
+-void
++__attribute__((target("s-memrealtime"))) void
+ OCKL_MANGLE_T(rtcwait,u32)(uint ticks)
+ {
+-    ulong now = __ockl_steadyctr_u64();
++    ulong now = __builtin_amdgcn_s_memrealtime();
+     ulong end = now + __builtin_amdgcn_readfirstlane(ticks);
+ 
+     if (__oclc_ISA_version >= 9000) {
+         while (end > now + 1625) {
+             __builtin_amdgcn_s_sleep(127);
+-            now = __ockl_steadyctr_u64();
++            now = __builtin_amdgcn_s_memrealtime();
+         }
+ 
+         while (end > now + 806) {
+             __builtin_amdgcn_s_sleep(63);
+-            now = __ockl_steadyctr_u64();
++            now = __builtin_amdgcn_s_memrealtime();
+         }
+ 
+         while (end > now + 396) {
+             __builtin_amdgcn_s_sleep(31);
+-            now = __ockl_steadyctr_u64();
++            now = __builtin_amdgcn_s_memrealtime();
+         }
+     }
+ 
+     while (end > now + 192) {
+         __builtin_amdgcn_s_sleep(15);
+-        now = __ockl_steadyctr_u64();
++        now = __builtin_amdgcn_s_memrealtime();
+     }
+ 
+     while (end > now + 89) {
+         __builtin_amdgcn_s_sleep(7);
+-        now = __ockl_steadyctr_u64();
++        now = __builtin_amdgcn_s_memrealtime();
+     }
+ 
+     while (end > now + 38) {
+         __builtin_amdgcn_s_sleep(3);
+-        now = __ockl_steadyctr_u64();
++        now = __builtin_amdgcn_s_memrealtime();
+     }
+ 
+     while (end > now) {
+         __builtin_amdgcn_s_sleep(1);
+-        now = __ockl_steadyctr_u64();
++        now = __builtin_amdgcn_s_memrealtime();
+     }
+ }
+ 
+-- 
+2.34.1
+
diff --git a/dev-libs/rocm-device-libs/rocm-device-libs-5.4.3.ebuild b/dev-libs/rocm-device-libs/rocm-device-libs-5.4.3.ebuild
new file mode 100644
index 000000000000..cab37f4bebff
--- /dev/null
+++ b/dev-libs/rocm-device-libs/rocm-device-libs-5.4.3.ebuild
@@ -0,0 +1,48 @@
+# Copyright 1999-2023 Gentoo Authors
+# Distributed under the terms of the GNU General Public License v2
+
+EAPI=8
+
+inherit cmake llvm
+
+LLVM_MAX_SLOT=15
+
+if [[ ${PV} == *9999 ]] ; then
+	EGIT_REPO_URI="https://github.com/RadeonOpenCompute/ROCm-Device-Libs/"
+	inherit git-r3
+	S="${WORKDIR}/${P}/src"
+else
+	SRC_URI="https://github.com/RadeonOpenCompute/ROCm-Device-Libs/archive/rocm-${PV}.tar.gz -> ${P}.tar.gz"
+	S="${WORKDIR}/ROCm-Device-Libs-rocm-${PV}"
+	KEYWORDS="~amd64"
+fi
+
+DESCRIPTION="Radeon Open Compute Device Libraries"
+HOMEPAGE="https://github.com/RadeonOpenCompute/ROCm-Device-Libs"
+
+LICENSE="MIT"
+SLOT="0/$(ver_cut 1-2)"
+IUSE="test"
+RESTRICT="!test? ( test )"
+
+RDEPEND="sys-devel/clang:${LLVM_MAX_SLOT}"
+DEPEND="${RDEPEND}"
+
+CMAKE_BUILD_TYPE=Release
+
+PATCHES=( "${FILESDIR}/${PN}-5.1.3-test-bitcode-dir.patch"
+	"${FILESDIR}/${PN}-5.1.3-llvm-link.patch"
+	"${FILESDIR}/${PN}-5.4.3-Revert-Update-counters-for-gfx11.patch" )
+
+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 "${LLVM_MAX_SLOT}")"
+	)
+	cmake_src_configure
+}