File rocclr-add-long-variants-for-__ffsll.patch of Package rocclr
From 8245dd0bb9ec719932372e8349a610d74c8cb514 Mon Sep 17 00:00:00 2001
From: Tom Rix <trix@redhat.com>
Date: Sun, 18 Jun 2023 05:54:24 -0700
Subject: [PATCH] add long variants for __ffsll
When building blender there is this error
/usr/include/nanovdb/NanoVDB.h:1880:12: error: call to '__ffsll' is ambiguous
return __ffsll(v);
^~~~~~~
/usr/include/hip/amd_detail/amd_device_functions.h:70:39: note: candidate function
__device__ static inline unsigned int __ffsll(unsigned long long int input) {
^
/usr/include/hip/amd_detail/amd_device_functions.h:78:39: note: candidate function
__device__ static inline unsigned int __ffsll(long long int input) {
on a 64 bit arch uint64_t can be an unsigned long, so add the long variants
Signed-off-by: Tom Rix <trix@redhat.com>
---
hipamd/include/hip/amd_detail/amd_device_functions.h | 8 ++++++++
1 file changed, 8 insertions(+)
diff --git a/hipamd/include/hip/amd_detail/amd_device_functions.h b/hipamd/include/hip/amd_detail/amd_device_functions.h
index ef9f127dd..5890d84cd 100644
--- a/hipamd/include/hip/amd_detail/amd_device_functions.h
+++ b/hipamd/include/hip/amd_detail/amd_device_functions.h
@@ -71,6 +71,10 @@ __device__ static inline unsigned int __ffsll(unsigned long long int input) {
return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1;
}
+__device__ static inline unsigned int __ffsll(unsigned long int input) {
+ return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1;
+}
+
__device__ static inline unsigned int __ffs(int input) {
return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1;
}
@@ -79,6 +83,10 @@ __device__ static inline unsigned int __ffsll(long long int input) {
return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1;
}
+__device__ static inline unsigned int __ffsll(long int input) {
+ return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1;
+}
+
// Given a 32/64-bit value exec mask and an integer value base (between 0 and WAVEFRONT_SIZE),
// find the n-th (given by offset) set bit in the exec mask from the base bit, and return the bit position.
// If not found, return -1.
--
2.41.0