From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from lists.gentoo.org (pigeon.gentoo.org [208.92.234.80]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (2048 bits)) (No client certificate requested) by finch.gentoo.org (Postfix) with ESMTPS id 180DD15800F for ; Wed, 22 Feb 2023 16:16:58 +0000 (UTC) Received: from pigeon.gentoo.org (localhost [127.0.0.1]) by pigeon.gentoo.org (Postfix) with SMTP id 61F4EE07D1; Wed, 22 Feb 2023 16:16:57 +0000 (UTC) Received: from smtp.gentoo.org (woodpecker.gentoo.org [140.211.166.183]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits)) (No client certificate requested) by pigeon.gentoo.org (Postfix) with ESMTPS id 2D423E07D3 for ; Wed, 22 Feb 2023 16:16:57 +0000 (UTC) Received: from oystercatcher.gentoo.org (oystercatcher.gentoo.org [148.251.78.52]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits)) (No client certificate requested) by smtp.gentoo.org (Postfix) with ESMTPS id 1D5C033BF51 for ; Wed, 22 Feb 2023 16:16:56 +0000 (UTC) Received: from localhost.localdomain (localhost [IPv6:::1]) by oystercatcher.gentoo.org (Postfix) with ESMTP id 89CCA8C1 for ; Wed, 22 Feb 2023 16:16:53 +0000 (UTC) From: "Sam James" To: gentoo-commits@lists.gentoo.org Content-Transfer-Encoding: 8bit Content-type: text/plain; charset=UTF-8 Reply-To: gentoo-dev@lists.gentoo.org, "Sam James" Message-ID: <1677082496.1adcab955fdbd4f3de4ff9f3a0900ebe8ce06fa3.sam@gentoo> Subject: [gentoo-commits] repo/gentoo:master commit in: dev-libs/rocm-device-libs/, dev-libs/rocm-device-libs/files/ X-VCS-Repository: repo/gentoo X-VCS-Files: dev-libs/rocm-device-libs/Manifest dev-libs/rocm-device-libs/files/rocm-device-libs-5.4.3-Revert-Update-counters-for-gfx11.patch dev-libs/rocm-device-libs/rocm-device-libs-5.4.3.ebuild X-VCS-Directories: dev-libs/rocm-device-libs/ dev-libs/rocm-device-libs/files/ X-VCS-Committer: sam X-VCS-Committer-Name: Sam James X-VCS-Revision: 1adcab955fdbd4f3de4ff9f3a0900ebe8ce06fa3 X-VCS-Branch: master Date: Wed, 22 Feb 2023 16:16:53 +0000 (UTC) Precedence: bulk List-Post: List-Help: List-Unsubscribe: List-Subscribe: List-Id: Gentoo Linux mail X-BeenThere: gentoo-commits@lists.gentoo.org X-Auto-Response-Suppress: DR, RN, NRN, OOF, AutoReply X-Archives-Salt: 66fe142c-2cd8-4b7c-997b-af5ffed3000b X-Archives-Hash: 1521d4c19a25a594b51ad5a45bc6e143 commit: 1adcab955fdbd4f3de4ff9f3a0900ebe8ce06fa3 Author: Niccolò Belli linuxsystems it> AuthorDate: Mon Feb 20 15:29:57 2023 +0000 Commit: Sam James gentoo org> CommitDate: Wed Feb 22 16:14:56 2023 +0000 URL: https://gitweb.gentoo.org/repo/gentoo.git/commit/?id=1adcab95 dev-libs/rocm-device-libs: add 5.4.3 Signed-off-by: Niccolò Belli linuxsystems.it> Signed-off-by: Sam James gentoo.org> dev-libs/rocm-device-libs/Manifest | 1 + ...bs-5.4.3-Revert-Update-counters-for-gfx11.patch | 216 +++++++++++++++++++++ .../rocm-device-libs/rocm-device-libs-5.4.3.ebuild | 48 +++++ 3 files changed, 265 insertions(+) diff --git a/dev-libs/rocm-device-libs/Manifest b/dev-libs/rocm-device-libs/Manifest index cf3568d4925e..2b7764830070 100644 --- a/dev-libs/rocm-device-libs/Manifest +++ b/dev-libs/rocm-device-libs/Manifest @@ -1,2 +1,3 @@ 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-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 +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 +}