From d1bbf23181e64a61a3be33f3b65f89d9c19626e3 Mon Sep 17 00:00:00 2001 From: Konstantin Pyzhov Date: Thu, 18 Apr 2019 13:07:24 -0400 Subject: [PATCH 1/2] Fix for __ffsll() device functions. [ROCm/clr commit: 5664ed320688f075f015946940135cfb03acb586] --- projects/clr/hipamd/include/hip/hcc_detail/device_functions.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index b88df2e168..7cb9406856 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h @@ -58,7 +58,7 @@ __device__ static inline unsigned int __ffs(unsigned int input) { } __device__ static inline unsigned int __ffsll(unsigned long long int input) { - return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1; + return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1; } __device__ static inline unsigned int __ffs(int input) { @@ -66,7 +66,7 @@ __device__ static inline unsigned int __ffs(int input) { } __device__ static inline unsigned int __ffsll(long long int input) { - return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1; + return ( input == 0 ? -1 : __builtin_ctzll(input) ) + 1; } __device__ static inline unsigned int __brev(unsigned int input) { From a525cc8f47aca6ab321f342c6b40b3a433c6f7f8 Mon Sep 17 00:00:00 2001 From: Konstantin Pyzhov Date: Fri, 19 Apr 2019 08:53:22 -0400 Subject: [PATCH 2/2] Fix for __popcll() device function implementation. [ROCm/clr commit: f6fbf8751dde7d3941f0886f61f249345c579cf7] --- projects/clr/hipamd/include/hip/hcc_detail/device_functions.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index 7cb9406856..044ca1e634 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h @@ -42,7 +42,7 @@ __device__ static inline unsigned int __popc(unsigned int input) { return __builtin_popcount(input); } __device__ static inline unsigned int __popcll(unsigned long long int input) { - return __builtin_popcountl(input); + return __builtin_popcountll(input); } __device__ static inline int __clz(int input) {